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

[lildmh]

Summary

This patch implements the parsing and sema support for OpenMP map clauses with potential user-defined mapper attached. User defined mapper is a new feature in OpenMP 5.0. A map clause can have an explicit or implicit associated mapper, which instructs the compiler to generate extra data mapping. 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 map(mapper(id) tofrom: ss) // use the mapper with name 'id' to map ss

[ABataev]

Also, need some kind of the stubbed codegen for the mapper

include/clang/AST/OpenMPClause.h

4236

The function has more than 10 parameters, it is not too good. Would be good to pack some of them into some kind of structure.

include/clang/Sema/Sema.h

9347

Also would be good to pack some of the params into structure

9431

Same, too many params

lib/Sema/SemaLookup.cpp

1074

Why do we need special processing of the mapper here?

1793

Again, what's so special in mapper lookup?

lib/Sema/SemaOpenMP.cpp

13187–13193

I think this can be simplified:

if (ER.isInvalid())
  continue;
MVLI.UDMapperList.push_back(ER.get());

13233–13239

Also can be simplified.

13364–13370

Again, simplify

lib/Sema/TreeTransform.h

8845

Enclose into braces

lib/Serialization/ASTReader.cpp

12312

i->I

[lildmh]

Hi Alexey,

Thanks very much for your quick review!

For the codegen, currently I'm thinking to do it within declare target codegen functions. So there is not a very clear interface to do mapper codegen (also, there is not a clear interface to do map clause codegen now), and that's why I didn't have a stub left in this patch. Please see other detailed responses inline.

include/clang/AST/OpenMPClause.h

4236

I can define a structure OMPMapClauseModifier within OMPMapClause to include all modifier related info. What do you think?

lib/Sema/SemaLookup.cpp

1074

The declare mapper lookup needs to find all OMPDeclareMapperDecls with the same name (the same as declare reduction lookup), and thus it goes through all scopes from the innermost to the outtermost.

Then it looks up a parent scope S (e.g., the outter {} in the following example), all OMPDeclareMapperDecls declared in the children scopes of S will appear in the range between IdResolver.begin(Name) and IdResolver.end(). Also, the declare mapper(id: struct S s) will appear before omp declare mapper(id: struct SS ss). This can cause the lookup function to ignore later omp declare mapper(id: struct SS ss) declared in the outter scope. As a result, we may not find the corrent mapper.

{
  #pragma omp declare mapper(id: struct S s) (s.a)
  {
    #pragma omp declare mapper(id: struct SS ss) (ss.b)
    ...
  }
}

To fix this problem, the purpose of this part of code is to ignore all OMPDeclareMapperDecls in inner scopes, which are already found in previous calls to LookupParsedName() from buildUserDefinedMapperRef.

I also found that the declare reduction lookup has the same problem. I'll push out a similar fix for the declare reduction later.

1793

The same as above. This is to fix C lookup, and the above is for C++

lib/Serialization/ASTReader.cpp

12312

I fixed it. Lower case i is also used in other loops in this function. Would you like me to fix them as well?

[lildmh]

Fix part of Alexey's comments

[ABataev]

In D58074#1394993, @lildmh wrote:

Hi Alexey,

Thanks very much for your quick review!

For the codegen, currently I'm thinking to do it within declare target codegen functions. So there is not a very clear interface to do mapper codegen (also, there is not a clear interface to do map clause codegen now), and that's why I didn't have a stub left in this patch. Please see other detailed responses inline.

This till must be handled somehow in codegen, better to ignore for now, I think.

include/clang/AST/OpenMPClause.h

4236

Yes, it would be good

lib/Sema/SemaLookup.cpp

1074

I don't think this is the correct behavior. For user-defined reductions, the standard says "This match is done by means of a name lookup in the base language." It means we should use the lookup rules of C/C++. For mapper, seems to me, we also should completely rely on the visibility/accessibility rules used by C/C++.

lib/Serialization/ASTReader.cpp

12312

This is the old code that was committed before the rules were defined. Soon they're going again to update the coding standard to allow lowCamel for the variables, but currently, only CamelCase is used for variables.
No, don't do this.

[lildmh]

include/clang/AST/OpenMPClause.h

4236

Then I'll do that. Thanks

lib/Sema/SemaLookup.cpp

1074

Again, thanks for your quick response!

{
  #pragma omp declare mapper(id: struct S s) (s.a)
  {
    #pragma omp declare mapper(id: struct SS ss) (ss.b)
    struct S kk;
    #pragma omp target map(mapper(id), tofrom: kk)
    {}
  }
}

If I don't add this code, the above code will report error that it cannnot find a mapper with name id for type struct S, because it can only find the second mapper for type struct SS. This doesn't look like correct behavior to me.

I think we are still following the lookup rules of C/C++ with this fix. It's because for declare mapper and reduction, we call LookupParsedName multiple times on different scopes. In other cases, LookupParsedName is always called once. Because of this difference, I think this fix is appropriate. What is your thought?

[ABataev]

lib/Sema/SemaLookup.cpp

1074

No, in your case we don't. According to the C/C++ rules, if the variable is redeclared, the latest declaration is used. The same rule must be applied to the mapper (according to the standard "The visibility and accessibility of this declaration are the same as those of a variable declared at the same point in the program.")
I think that the code should fail in this case, because you would the error message in case of the variables declarations with the same names in those scopes.

[lildmh]

lib/Sema/SemaLookup.cpp

1074

Hi Alexey,

Thanks for your quick response! I don't think it's redeclared in this case, because a mapper has two properties: name and type, just as declare reduction. Only when both name and type are the same, it should be considered as redeclaration.

If we consider the above example as redeclaration of a mapper, then we can always just call LookupParsedName once to find the most relevant mapper/reductor. Then why do you need to search reductors in a while loop in buildDeclareReductionRef?

Also, the following example will be correct using the original lookup functions, without my fix:

{
  #pragma omp declare mapper(id: struct S s) (s.a)
    #pragma omp declare mapper(id: struct SS ss) (ss.b)
    struct S kk;
    #pragma omp target map(mapper(id), tofrom: kk)
    {}
}

If we consider the second mapper as a redeclaration of the first one, this should fail as well. What do you think?

[ABataev]

lib/Sema/SemaLookup.cpp

1074

The standard clearly states that "The visibility and accessibility of this declaration are the same as those of a variable declared at the same point in the program.". For variables (btw, they also have 2 attributes - name and type) the inner declaration makes the outer one not visible. The same rule should be applied to the mapper.
Instead of this case, I'd suggest to implement the argument-dependent lookup for the declarations with the same name but with the different types, declared in the same scope.
Your second example should not produce the error as the mappers are declared in the same scope.

[lildmh]

lib/Sema/SemaLookup.cpp

1074

Thanks for your patience! I think I got your point about the visibility and accessibility rule. Now I'm confused about the implementation of declare reduction lookup. Could you explain why you use a while loop to call LookupParsedName in buildDeclareReductionRef? I thought I understood your intention when I read your code, but I'm confused now.

For argument dependent lookup, I thought it is for functions. Declare reduction can be considered as a function, but I don't think mappers should be considered as functions. What do you think?

[ABataev]

lib/Sema/SemaLookup.cpp

1074

It just scans through the scopes until it finds at least one UDR with the specified name, nothing else. It does not try to find all the UDRs with the given name in all scopes.
As for the ADL, mapper, in this case, is similar to UDR because you need to find the mapper with the required name and type.

[kkwli0]

include/clang/AST/OpenMPClause.h

4077

assoicated -> associated

4212

amount -> number?

4229

assoicated -> associated

lib/Parse/ParseOpenMP.cpp

2144

Although it is an error situation, will we have different behavior?

... map(xclose, to: x)

Previously, it always parses both modifier and type. After the change, the call of parseMapType is skipped. If it OMPC_MAP_unknown, IsMapTypeImplicit is set.

[lildmh]

lib/Parse/ParseOpenMP.cpp

2144

Hi Kelvin,

Thanks a lot for the review! It will not change the behavior of parsing modifiers other than mapper. parseMapTypeModifiers will only return true when it finds a invalid mapper for now.

[lildmh]

1. Move mapper related info to OMPMappableExprListClause, so that to and from clauses can utilize the same infrastructure as well (real support for them will be in another patch);
2. Combine the function parameters for declare mapper and declare reduction, to reduce the function parameter number of ActOnOMPVarListClause;
3. Implement ADL for mapper, and remove mapper lookups in SemaLookup.cpp;
4. Didn't combine modifier related info into a structure, because for to and from clauses will only have possible mapper modifier, not close and always;
5. Address typos that Kelvin found;
6. Rebase.

[ABataev]

Check, please, that adding this to the map clause does not crash the codegen. Would be good to ignore this construct in codegen for now, if it is used in the code.

include/clang/AST/OpenMPClause.h

4116

Also would be good to pack the parameters into the structure.

[lildmh]

Hi Alexey,

Again thanks for your review! The codegen completely ignores any mapper related info for now, so it should not crash map clause codegen. It also passed the regression test, so map clause codegen should be fine.

include/clang/AST/OpenMPClause.h

4116

Among these parameters:

1. 2 of them are for modifiers. Since mapper can also be used in to and from clauses, it is not good to combine mapper info with them.
2. 2 of them are mapper info. I can combine MapperQualifierLoc and MapperIdInfo into a structure. It's not a big saving though (only reduce one parameter). Do you want me to do that?
3. 4 of them are number of vars... It seems not a good idea to combine them into one structure.
4. The rest are locations. They are kinda all over the place. It doesn't seem a good idea to combine them as well.

[ABataev]

In D58074#1398187, @lildmh wrote:

Hi Alexey,

Again thanks for your review! The codegen completely ignores any mapper related info for now, so it should not crash map clause codegen. It also passed the regression test, so map clause codegen should be fine.

Hi Lingda, no problems!
It would be good to have at least one codegen test with the mapper construct to prove that it does not crash the codegen. We don't have any codegen test for this new construct.

include/clang/AST/OpenMPClause.h

4116

1. The same structure must be used for to and from clauses, not a problem.

2-4. you can combine it with the number of vars and locations into one structure, no?

[lildmh]

Sure I'll add a codegen test with mapper. Thanks!

include/clang/AST/OpenMPClause.h

4116

None of these locations are for mapper, so I think it's not good to combine them with mapper.
Those numbers (NumVars, NumUniqueDeclarations, NumComponentLists, NumComponents) are not for mapper either. NumVars can be viewed as partially for mapper, but its main purpose is for the number of list items in map clause. So logically, I think they should not be combined with mapper.
What do you think?

[ABataev]

include/clang/AST/OpenMPClause.h

4116

It is not about mapper. We just have too many params in functions and need to reduce it somehow. I don't care about the logic of the parameters packing.

[lildmh]

Introduce a structure OMPMappableExprListSizeTy within OMPMappableExprListClause to aggregate all 4 sizes needed by such clause, and thus reduce the number of parameters when creating a map clause. This can also be used by other similar clauses. I'll have another patch to let other clauses to use this structure as well.

[ABataev]

include/clang/AST/OpenMPClause.h

3626

Add default initializers for the fields

3639

Remove this constructor, we should use only the one with OMPMappableExprListSizeTy

3661–3664

I think it is worth to add SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc, to OMPMappableExprListSizeTy. Of course, you need to rename the structure, something like OMPMappableClauseData is good enough. Or you can pack them into a different structure.

4237

Also too many params, need to "squash" some of them into a structure, maybe several structures.

include/clang/Sema/Sema.h

9433

Also too many params here.

[lildmh]

Further reduce the number of parameters for map, to, from, use_device, and is_device clauses. Also rebase.

include/clang/AST/OpenMPClause.h

3661–3664

I put StartLoc, LParenLoc, EndLoc into another struct OMPMappableExprListLocTy. I didn't combine them with OMPMappableExprListSizeTy, because some fuctions only use one of them, and some other functions use both of them.

Also modify to, from, is_device, and use_device clauses to use the same interfaces.

[ABataev]

include/clang/AST/OpenMPClause.h

3649

Pass those new structures by const reference, not by value. Do this for all the functions.

[lildmh]

Change parameters to const &

[ABataev]

include/clang/Sema/Sema.h

9347

What about this function?

[lildmh]

Thanks for the catch! I also change the name from OMPMappableExprListLocTy to OMPVarListLocTy to be more accurate. We can potentially factorize other varlist clause code with this. But it will be too large for this patch.

[ABataev]

In D58074#1399881, @lildmh wrote:

Thanks for the catch! I also change the name from OMPMappableExprListLocTy to OMPVarListLocTy to be more accurate. We can potentially factorize other varlist clause code with this. But it will be too large for this patch.

Thanks for fixing this. Yes, if you can think we improve something in the code, please, go ahead in the next patches. LGTM!

[lildmh]

Rebase

include/clang/AST/OpenMPClause.h

protected:
   /// NumComponentLists total component lists, and \a NumComponents total
   /// components.
   ///
   /// \param K Kind of the clause.
   /// \param StartLoc Starting location of the clause (the clause keyword).
   /// \param LParenLoc Location of '('.
   /// \param EndLoc Ending location of the clause.
   /// \param NumVars Number of expressions listed in the clause.
   /// \param NumUniqueDeclarations Number of unique base declarations in this
   /// clause.
   /// \param NumComponentLists Number of component lists in this clause - one
   /// list for each expression in the clause.
   /// \param NumComponents Total number of expression components in the clause.
   OMPMappableExprListClause(OpenMPClauseKind K, SourceLocation StartLoc,
                             SourceLocation LParenLoc, SourceLocation EndLoc,
                             unsigned NumVars, unsigned NumUniqueDeclarations,
                             unsigned NumComponentLists, unsigned NumComponents)
       : OMPVarListClause<T>(K, StartLoc, LParenLoc, EndLoc, NumVars),
         NumUniqueDeclarations(NumUniqueDeclarations),
         NumComponentLists(NumComponentLists), NumComponents(NumComponents) {}
 
   /// Get the unique declarations that are in the trailing objects of the
   /// class.
   MutableArrayRef<ValueDecl *> getUniqueDeclsRef() {
     return MutableArrayRef<ValueDecl *>(
         static_cast<T *>(this)->template getTrailingObjects<ValueDecl *>(),
         NumUniqueDeclarations);
   }
 
   /// Get the unique declarations that are in the trailing objects of the
   /// class.
   ArrayRef<ValueDecl *> getUniqueDeclsRef() const {
     return ArrayRef<ValueDecl *>(
         static_cast<const T *>(this)
             ->template getTrailingObjects<ValueDecl *>(),
         NumUniqueDeclarations);
   }
 
   /// Set the unique declarations that are in the trailing objects of the
   /// class.
   void setUniqueDecls(ArrayRef<ValueDecl *> UDs) {
     assert(UDs.size() == NumUniqueDeclarations &&
            "Unexpected amount of unique declarations.");
     std::copy(UDs.begin(), UDs.end(), getUniqueDeclsRef().begin());
   }
 
   /// Get the number of lists per declaration that are in the trailing
   /// objects of the class.
   MutableArrayRef<unsigned> getDeclNumListsRef() {
     return MutableArrayRef<unsigned>(
         static_cast<T *>(this)->template getTrailingObjects<unsigned>(),
         NumUniqueDeclarations);
   }
 
   /// Get the number of lists per declaration that are in the trailing
class OMPMapClause final : public OMPMappableExprListClause<OMPMapClause>,
   friend class OMPClauseReader;
   friend OMPMappableExprListClause;
   friend OMPVarListClause;
   friend TrailingObjects;
 
   /// 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();
+    return 2 * varlist_size();
   }
   size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
     return getUniqueDeclarationsNum();
   }
   size_t numTrailingObjects(OverloadToken<unsigned>) const {
     return getUniqueDeclarationsNum() + getTotalComponentListNum();
   }
 
 public:
   /// Number of allowed map-type-modifiers.
   static constexpr unsigned NumberOfModifiers =
       OMPC_MAP_MODIFIER_last - OMPC_MAP_MODIFIER_unknown - 1;
 
 private:
   /// Map-type-modifiers for the 'map' clause.
   OpenMPMapModifierKind MapTypeModifiers[NumberOfModifiers] = {
-    OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown
-  };
-  
+      OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown,
+      OMPC_MAP_MODIFIER_unknown};
   /// Location of map-type-modifiers for the 'map' clause.
   SourceLocation MapTypeModifiersLoc[NumberOfModifiers];
 
+  /// C++ nested name specifier for the assoicated user-defined mapper.
+  NestedNameSpecifierLoc MapperQualifierLoc;
+
+  /// The associated user-defined mapper identifier information.
+  DeclarationNameInfo MapperIdInfo;
+
   /// Map type for the 'map' clause.
   OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
 
   /// Is this an implicit map type or not.
   bool MapTypeIsImplicit = false;
 
   /// Location of the map type.
   SourceLocation MapLoc;
 
   /// Colon location.
   SourceLocation ColonLoc;
 
   /// Build a clause for \a NumVars listed expressions, \a
   /// NumUniqueDeclarations declarations, \a NumComponentLists total component
   /// lists, and \a NumComponents total expression components.
   ///
   /// \param MapModifiers Map-type-modifiers.
   /// \param MapModifiersLoc Locations of map-type-modifiers.
+  /// \param MapperQualifierLoc C++ nested name specifier for the assoicated
+  /// user-defined mapper.
+  /// \param MapperIdInfo The identifier of associated user-defined mapper.
   /// \param MapType Map type.
   /// \param MapTypeIsImplicit Map type is inferred implicitly.
   /// \param MapLoc Location of the map type.
   /// \param StartLoc Starting location of the clause.
   /// \param EndLoc Ending location of the clause.
   /// \param NumVars Number of expressions listed in this clause.
   /// \param NumUniqueDeclarations Number of unique base declarations in this
   /// clause.
   /// \param NumComponentLists Number of component lists in this clause.
   /// \param NumComponents Total number of expression components in the clause.
   explicit OMPMapClause(ArrayRef<OpenMPMapModifierKind> MapModifiers,
                         ArrayRef<SourceLocation> MapModifiersLoc,
+                        NestedNameSpecifierLoc MapperQualifierLoc,
+                        DeclarationNameInfo MapperIdInfo,
                         OpenMPMapClauseKind MapType, bool MapTypeIsImplicit,
                         SourceLocation MapLoc, SourceLocation StartLoc,
                         SourceLocation LParenLoc, SourceLocation EndLoc,
                         unsigned NumVars, unsigned NumUniqueDeclarations,
                         unsigned NumComponentLists, unsigned NumComponents)
       : OMPMappableExprListClause(OMPC_map, StartLoc, LParenLoc, EndLoc,
                                   NumVars, NumUniqueDeclarations,
                                   NumComponentLists, NumComponents),
-        MapType(MapType), MapTypeIsImplicit(MapTypeIsImplicit),
-        MapLoc(MapLoc) {
-          assert(llvm::array_lengthof(MapTypeModifiers) == MapModifiers.size()
-                 && "Unexpected number of map type modifiers.");
-          llvm::copy(MapModifiers, std::begin(MapTypeModifiers));
+        MapperQualifierLoc(MapperQualifierLoc), MapperIdInfo(MapperIdInfo),
+        MapType(MapType), MapTypeIsImplicit(MapTypeIsImplicit), MapLoc(MapLoc) {
+    assert(llvm::array_lengthof(MapTypeModifiers) == MapModifiers.size() &&
+           "Unexpected number of map type modifiers.");
 
-          assert(llvm::array_lengthof(MapTypeModifiersLoc) ==
-                     MapModifiersLoc.size() &&
-                 "Unexpected number of map type modifier locations.");
-          llvm::copy(MapModifiersLoc, std::begin(MapTypeModifiersLoc));
   }
 
   /// Build an empty clause.
   ///
   /// \param NumVars Number of expressions listed in this clause.
   /// \param NumUniqueDeclarations Number of unique base declarations in this
   /// clause.
   /// \param NumComponentLists Number of component lists in this clause.
private:
   /// Set type location.
   ///
   /// \param TLoc Type location.
   void setMapLoc(SourceLocation TLoc) { MapLoc = TLoc; }
 
   /// Set colon location.
   void setColonLoc(SourceLocation Loc) { ColonLoc = Loc; }
 
+  /// Set the nested name specifier of associated user-defined mapper.
+  void setMapperQualifierLoc(NestedNameSpecifierLoc NNSL) {
+    MapperQualifierLoc = NNSL;
+  }
+
+  /// Set the name of associated user-defined mapper.
+  void setMapperIdInfo(DeclarationNameInfo MapperId) {
+    MapperIdInfo = MapperId;
+  }
+
+  /// Get the user-defined mapper references that are in the trailing objects of
+  /// the class.
+  MutableArrayRef<Expr *> getUDMapperRefs() {
+    return llvm::makeMutableArrayRef<Expr *>(
+        getTrailingObjects<Expr *>() + varlist_size(), varlist_size());
+  }
+
+  /// Get the user-defined mappers references that are in the trailing objects
+  /// of the class.
+  ArrayRef<Expr *> getUDMapperRefs() const {
+    return llvm::makeArrayRef<Expr *>(
+        getTrailingObjects<Expr *>() + varlist_size(), varlist_size());
+  }
+
+  /// Set the user-defined mappers that are in the trailing objects of the
+  /// class.
+  void setUDMapperRefs(ArrayRef<Expr *> DMDs) {
+    assert(DMDs.size() == varlist_size() &&
+           "Unexpected amount of user-defined mappers.");
+    std::copy(DMDs.begin(), DMDs.end(), getUDMapperRefs().begin());
+  }
+
 public:
   /// Creates clause with a list of variables \a VL.
   ///
   /// \param C AST context.
   /// \param StartLoc Starting location of the clause.
   /// \param 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 MapModifiers Map-type-modifiers.
   /// \param MapModifiersLoc Location of map-type-modifiers.
+  /// \param UDMQualifierLoc C++ nested name specifier for the assoicated
+  /// user-defined mapper.
+  /// \param MapperId The identifier of associated user-defined mapper.
   /// \param Type Map type.
   /// \param TypeIsImplicit Map type is inferred implicitly.
   /// \param TypeLoc Location of the map type.
-  static OMPMapClause *Create(const ASTContext &C, SourceLocation StartLoc,
-                              SourceLocation LParenLoc, SourceLocation EndLoc,
-                              ArrayRef<Expr *> Vars,
-                              ArrayRef<ValueDecl *> Declarations,
-                              MappableExprComponentListsRef ComponentLists,
-                              ArrayRef<OpenMPMapModifierKind> MapModifiers,
-                              ArrayRef<SourceLocation> MapModifiersLoc,
-                              OpenMPMapClauseKind Type, bool TypeIsImplicit,
-                              SourceLocation TypeLoc);
+  static OMPMapClause *
+  Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation LParenLoc,
+         SourceLocation EndLoc, ArrayRef<Expr *> Vars,
+         ArrayRef<Expr *> UDMapperRefs,
+         NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId,
+         OpenMPMapClauseKind Type, bool TypeIsImplicit, SourceLocation TypeLoc);
 
   /// Creates an empty clause with the place for \a NumVars original
   /// expressions, \a NumUniqueDeclarations declarations, \NumComponentLists
   /// lists, and \a NumComponents expression components.
   ///
   /// \param C AST context.
   /// \param NumVars Number of expressions listed in the clause.
   /// \param NumUniqueDeclarations Number of unique base declarations in this
OpenMPMapModifierKind getMapTypeModifier(unsigned Cnt) const LLVM_READONLY {
     assert(Cnt < NumberOfModifiers &&
            "Requested modifier exceeds the total number of modifiers.");
     return MapTypeModifiers[Cnt];
   }
 
   /// Fetches the map-type-modifier location at 'Cnt' index of array of
   /// modifiers' locations.
   ///
-  /// \param Cnt index for map-type-modifier location.  
+  /// \param Cnt index for map-type-modifier location.
   SourceLocation getMapTypeModifierLoc(unsigned Cnt) const LLVM_READONLY {
     assert(Cnt < NumberOfModifiers &&
            "Requested modifier location exceeds total number of modifiers.");
     return MapTypeModifiersLoc[Cnt];
   }
 
   /// Fetches ArrayRef of map-type-modifiers.
   ArrayRef<OpenMPMapModifierKind> getMapTypeModifiers() const LLVM_READONLY {
     return llvm::makeArrayRef(MapTypeModifiers);
   }
-  
   /// Fetches ArrayRef of location of map-type-modifiers.
   ArrayRef<SourceLocation> getMapTypeModifiersLoc() const LLVM_READONLY {
     return llvm::makeArrayRef(MapTypeModifiersLoc);
   }
 
+  /// Gets the name info for associated user-defined mapper.
+  const DeclarationNameInfo &getMapperIdInfo() const { return MapperIdInfo; }
+
+  /// Gets the nested name specifier for associated user-defined mapper.
+  NestedNameSpecifierLoc getMapperQualifierLoc() const {
+    return MapperQualifierLoc;
+  }
+
   /// Fetches location of clause mapping kind.
   SourceLocation getMapLoc() const LLVM_READONLY { return MapLoc; }
 
   /// Get colon location.
   SourceLocation getColonLoc() const { return ColonLoc; }
 
+  using mapperlist_iterator = MutableArrayRef<Expr *>::iterator;
+  using mapperlist_const_iterator = ArrayRef<const Expr *>::iterator;
+  using mapperlist_range = llvm::iterator_range<mapperlist_iterator>;
+  using mapperlist_const_range =
+      llvm::iterator_range<mapperlist_const_iterator>;
+
+  mapperlist_iterator mapperlist_begin() { return getUDMapperRefs().begin(); }
+  mapperlist_iterator mapperlist_end() { return getUDMapperRefs().end(); }
+  mapperlist_const_iterator mapperlist_begin() const {
+    return getUDMapperRefs().begin();
+  }
+  mapperlist_const_iterator mapperlist_end() const {
+    return getUDMapperRefs().end();
+  }
+  mapperlist_range mapperlists() {
+    return mapperlist_range(mapperlist_begin(), mapperlist_end());
+  }
+  mapperlist_const_range mapperlists() const {
+    return mapperlist_const_range(mapperlist_begin(), mapperlist_end());
+  }
+
   child_range children() {
     return child_range(
         reinterpret_cast<Stmt **>(varlist_begin()),
         reinterpret_cast<Stmt **>(varlist_end()));
   }
 
   static bool classof(const OMPClause *T) {
     return T->getClauseKind() == OMPC_map;

include/clang/Basic/DiagnosticParseKinds.td

def err_omp_expected_identifier_for_critical : Error<
   "expected identifier specifying the name of the 'omp critical' directive">;
 def err_omp_expected_reduction_identifier : Error<
   "expected identifier or one of the following operators: '+', '-', '*', '&', '|', '^', '&&', or '||'">;
 def err_omp_decl_in_declare_simd : Error<
   "function declaration is expected after 'declare simd' directive">;
 def err_omp_unknown_map_type : Error<
   "incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'">;
 def err_omp_unknown_map_type_modifier : Error<
-  "incorrect map type modifier, expected 'always' or 'close'">;
+  "incorrect map type modifier, expected 'always', 'close', or 'mapper'">;
 def err_omp_map_type_missing : Error<
   "missing map type">;
 def err_omp_map_type_modifier_missing : Error<
   "missing map type modifier">;
 def err_omp_declare_simd_inbranch_notinbranch : Error<
   "unexpected '%0' clause, '%1' is specified already">;
 def err_expected_end_declare_target : Error<
   "expected '#pragma omp end declare target'">;

include/clang/Basic/DiagnosticSemaKinds.td

 def err_omp_wrong_var_in_declare_reduction : Error<"only %select{'omp_priv' or 'omp_orig'|'omp_in' or 'omp_out'}0 variables are allowed in %select{initializer|combiner}0 expression">;
 def err_omp_declare_reduction_redefinition : Error<"redefinition of user-defined reduction for type %0">;
 def err_omp_mapper_wrong_type : Error<
   "mapper type must be of struct, union or class type">;
 def err_omp_declare_mapper_wrong_var : Error<
   "only variable %0 is allowed in map clauses of this 'omp declare mapper' directive">;
 def err_omp_declare_mapper_redefinition : Error<
   "redefinition of user-defined mapper for type %0 with name %1">;
+def err_omp_invalid_mapper: Error<
+  "cannot find a valid user-defined mapper for type %0 with name %1">;
 def err_omp_array_section_use : Error<"OpenMP array section is not allowed here">;
 def err_omp_typecheck_section_value : Error<
   "subscripted value is not an array or pointer">;
 def err_omp_typecheck_section_not_integer : Error<
   "array section %select{lower bound|length}0 is not an integer">;
 def err_omp_section_function_type : Error<
   "section of pointer to function type %0">;
 def warn_omp_section_is_char : Warning<"array section %select{lower bound|length}0 is of type 'char'">,

include/clang/Basic/OpenMPKinds.def

 OPENMP_MAP_KIND(from)
 OPENMP_MAP_KIND(tofrom)
 OPENMP_MAP_KIND(delete)
 OPENMP_MAP_KIND(release)
 
 // Map-type-modifiers for 'map' clause.
 OPENMP_MAP_MODIFIER_KIND(always)
 OPENMP_MAP_MODIFIER_KIND(close)
+OPENMP_MAP_MODIFIER_KIND(mapper)
 
 // Clauses allowed for OpenMP directive 'taskloop'.
 OPENMP_TASKLOOP_CLAUSE(if)
 OPENMP_TASKLOOP_CLAUSE(shared)
 OPENMP_TASKLOOP_CLAUSE(private)
 OPENMP_TASKLOOP_CLAUSE(firstprivate)
 OPENMP_TASKLOOP_CLAUSE(lastprivate)
 OPENMP_TASKLOOP_CLAUSE(default)

include/clang/Parse/Parser.h

struct OpenMPVarListDataTy {
     CXXScopeSpec ReductionIdScopeSpec;
     DeclarationNameInfo ReductionId;
     OpenMPDependClauseKind DepKind = OMPC_DEPEND_unknown;
     OpenMPLinearClauseKind LinKind = OMPC_LINEAR_val;
     SmallVector<OpenMPMapModifierKind, OMPMapClause::NumberOfModifiers>
     MapTypeModifiers;
     SmallVector<SourceLocation, OMPMapClause::NumberOfModifiers>
     MapTypeModifiersLoc;
+    CXXScopeSpec MapperIdScopeSpec;
+    DeclarationNameInfo MapperId;
     OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
     bool IsMapTypeImplicit = false;
     SourceLocation DepLinMapLoc;
   };
 
   /// Parses clauses with list.
   bool ParseOpenMPVarList(OpenMPDirectiveKind DKind, OpenMPClauseKind Kind,
                           SmallVectorImpl<Expr *> &Vars,
                           OpenMPVarListDataTy &Data);
   bool ParseUnqualifiedId(CXXScopeSpec &SS, bool EnteringContext,
                           bool AllowDestructorName,
                           bool AllowConstructorName,
                           bool AllowDeductionGuide,
                           ParsedType ObjectType,
                           SourceLocation *TemplateKWLoc,
                           UnqualifiedId &Result);
+  /// Parses map-type-modifiers in map clause.
+  /// map([ [map-type-modifier[,] [map-type-modifier[,] ...] map-type : ] list)
+  /// where, map-type-modifier ::= always | close | mapper(mapper-identifier)
+  bool parseMapTypeModifiers(OpenMPVarListDataTy &Data);
 
 private:
   //===--------------------------------------------------------------------===//
   // C++ 14: Templates [temp]
 
   // C++ 14.1: Template Parameters [temp.param]
   Decl *ParseDeclarationStartingWithTemplate(DeclaratorContext Context,
                                              SourceLocation &DeclEnd,

include/clang/Sema/Sema.h

OMPClause *ActOnOpenMPVarListClause(
       OpenMPClauseKind Kind, ArrayRef<Expr *> Vars, Expr *TailExpr,
       SourceLocation StartLoc, SourceLocation LParenLoc,
       SourceLocation ColonLoc, SourceLocation EndLoc,
       CXXScopeSpec &ReductionIdScopeSpec,
       const DeclarationNameInfo &ReductionId, OpenMPDependClauseKind DepKind,
       OpenMPLinearClauseKind LinKind,
       ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
       ArrayRef<SourceLocation> MapTypeModifiersLoc,
+      CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId,
       OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
       SourceLocation DepLinMapLoc);
   /// Called on well-formed 'private' clause.
   OMPClause *ActOnOpenMPPrivateClause(ArrayRef<Expr *> VarList,
                                       SourceLocation StartLoc,
                                       SourceLocation LParenLoc,
                                       SourceLocation EndLoc);
   /// Called on well-formed 'firstprivate' clause.
ActOnOpenMPDependClause(OpenMPDependClauseKind DepKind, SourceLocation DepLoc,
                           SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
                           SourceLocation StartLoc, SourceLocation LParenLoc,
                           SourceLocation EndLoc);
   /// Called on well-formed 'device' clause.
   OMPClause *ActOnOpenMPDeviceClause(Expr *Device, SourceLocation StartLoc,
                                      SourceLocation LParenLoc,
                                      SourceLocation EndLoc);
   /// Called on well-formed 'map' clause.
-  OMPClause *
-  ActOnOpenMPMapClause(ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
-                       ArrayRef<SourceLocation> MapTypeModifiersLoc,
-                       OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
-                       SourceLocation MapLoc, SourceLocation ColonLoc,
-                       ArrayRef<Expr *> VarList, SourceLocation StartLoc,
-                       SourceLocation LParenLoc, SourceLocation EndLoc);
+  OMPClause *ActOnOpenMPMapClause(
+      ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
+      CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId,
+      SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
+      SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc,
+      ArrayRef<Expr *> UnresolvedMappers = llvm::None);
   /// Called on well-formed 'num_teams' clause.
   OMPClause *ActOnOpenMPNumTeamsClause(Expr *NumTeams, SourceLocation StartLoc,
                                        SourceLocation LParenLoc,
                                        SourceLocation EndLoc);
   /// Called on well-formed 'thread_limit' clause.
   OMPClause *ActOnOpenMPThreadLimitClause(Expr *ThreadLimit,
                                           SourceLocation StartLoc,
                                           SourceLocation LParenLoc,

lib/AST/OpenMPClause.cpp

for (const ValueDecl *D : Declarations) {
     if (Cache.count(VD))
       continue;
     ++TotalNum;
     Cache.insert(VD);
   }
   return TotalNum;
 }
 
-OMPMapClause *
-OMPMapClause::Create(const ASTContext &C, SourceLocation StartLoc,
-                     SourceLocation LParenLoc, SourceLocation EndLoc,
-                     ArrayRef<Expr *> Vars, ArrayRef<ValueDecl *> Declarations,
-                     MappableExprComponentListsRef ComponentLists,
-                     ArrayRef<OpenMPMapModifierKind> MapModifiers,
-                     ArrayRef<SourceLocation> MapModifiersLoc,
-                     OpenMPMapClauseKind Type, bool TypeIsImplicit,
-                     SourceLocation TypeLoc) {
+OMPMapClause *OMPMapClause::Create(
+    const ASTContext &C, SourceLocation StartLoc, SourceLocation LParenLoc,
+    SourceLocation EndLoc, ArrayRef<Expr *> Vars,
+    ArrayRef<ValueDecl *> Declarations,
+    MappableExprComponentListsRef ComponentLists, ArrayRef<Expr *> UDMapperRefs,
+    NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId,
+    OpenMPMapClauseKind Type, bool TypeIsImplicit, SourceLocation TypeLoc) {
   unsigned NumVars = Vars.size();
   unsigned NumUniqueDeclarations =
       getUniqueDeclarationsTotalNumber(Declarations);
   unsigned NumComponentLists = ComponentLists.size();
   unsigned NumComponents = getComponentsTotalNumber(ComponentLists);
 
   // We need to allocate:
-  // NumVars x Expr* - we have an original list expression for each clause list
-  // entry.
+  // 2 x NumVars x Expr* - we have an original list expression and an associated
+  // user-defined mapper for each clause list entry.
   // NumUniqueDeclarations x ValueDecl* - unique base declarations associated
   // with each component list.
   // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
   // 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>(
-          NumVars, NumUniqueDeclarations,
+          2 * NumVars, NumUniqueDeclarations,
           NumUniqueDeclarations + NumComponentLists, NumComponents));
   OMPMapClause *Clause = new (Mem) OMPMapClause(
-      MapModifiers, MapModifiersLoc, Type, TypeIsImplicit, TypeLoc, StartLoc,
-      LParenLoc, EndLoc, NumVars, NumUniqueDeclarations, NumComponentLists,
-      NumComponents);
+      MapModifiers, MapModifiersLoc, UDMQualifierLoc, MapperId, Type,
+      TypeIsImplicit, TypeLoc, StartLoc, LParenLoc, EndLoc, NumVars,
+      NumUniqueDeclarations, NumComponentLists, NumComponents);
 
   Clause->setVarRefs(Vars);
+  Clause->setUDMapperRefs(UDMapperRefs);
   Clause->setClauseInfo(Declarations, ComponentLists);
   Clause->setMapType(Type);
   Clause->setMapLoc(TypeLoc);
   return Clause;
 }
 
 OMPMapClause *OMPMapClause::CreateEmpty(const ASTContext &C, unsigned NumVars,
                                         unsigned NumUniqueDeclarations,
                                         unsigned NumComponentLists,
                                         unsigned NumComponents) {
   void *Mem = C.Allocate(
       totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
-          NumVars, NumUniqueDeclarations,
+          2 * NumVars, NumUniqueDeclarations,
           NumUniqueDeclarations + NumComponentLists, NumComponents));
   return new (Mem) OMPMapClause(NumVars, NumUniqueDeclarations,
                                 NumComponentLists, NumComponents);
 }
 
 OMPToClause *OMPToClause::Create(const ASTContext &C, SourceLocation StartLoc,
                                  SourceLocation LParenLoc,
                                  SourceLocation EndLoc, ArrayRef<Expr *> Vars,

 void OMPClausePrinter::VisitOMPMapClause(OMPMapClause *Node) {
   if (!Node->varlist_empty()) {
     OS << "map(";
     if (Node->getMapType() != OMPC_MAP_unknown) {
       for (unsigned I = 0; I < OMPMapClause::NumberOfModifiers; ++I) {
         if (Node->getMapTypeModifier(I) != OMPC_MAP_MODIFIER_unknown) {
           OS << getOpenMPSimpleClauseTypeName(OMPC_map,
                                               Node->getMapTypeModifier(I));
+          if (Node->getMapTypeModifier(I) == OMPC_MAP_MODIFIER_mapper) {
+            OS << '(';
+            NestedNameSpecifier *MapperNNS =
+                Node->getMapperQualifierLoc().getNestedNameSpecifier();
+            if (MapperNNS)
+              MapperNNS->print(OS, Policy);
+            OS << Node->getMapperIdInfo() << ')';
+          }
           OS << ',';
         }
       }
       OS << getOpenMPSimpleClauseTypeName(OMPC_map, Node->getMapType());
       OS << ':';
     }
     VisitOMPClauseList(Node, ' ');
     OS << ")";

lib/Parse/ParseOpenMP.cpp

static OpenMPMapModifierKind isMapModifier(Parser &P) {
   Preprocessor &PP = P.getPreprocessor();
   OpenMPMapModifierKind TypeModifier = static_cast<OpenMPMapModifierKind>(
       getOpenMPSimpleClauseType(OMPC_map, PP.getSpelling(Tok)));
   return TypeModifier;
 }
 
 /// Parse map-type-modifiers in map clause.
 /// map([ [map-type-modifier[,] [map-type-modifier[,] ...] map-type : ] list)
-/// where, map-type-modifier ::= always | close 
-static void parseMapTypeModifiers(Parser &P,
-                                  Parser::OpenMPVarListDataTy &Data) {
-  Preprocessor &PP = P.getPreprocessor();
-  while (P.getCurToken().isNot(tok::colon)) {
-    Token Tok = P.getCurToken();
-    OpenMPMapModifierKind TypeModifier = isMapModifier(P);
+/// where, map-type-modifier ::= always | close | mapper(mapper-identifier)
+bool Parser::parseMapTypeModifiers(OpenMPVarListDataTy &Data) {
+  while (getCurToken().isNot(tok::colon)) {
+    OpenMPMapModifierKind TypeModifier = isMapModifier(*this);
     if (TypeModifier == OMPC_MAP_MODIFIER_always ||
         TypeModifier == OMPC_MAP_MODIFIER_close) {
       Data.MapTypeModifiers.push_back(TypeModifier);
       Data.MapTypeModifiersLoc.push_back(Tok.getLocation());
-      P.ConsumeToken();
+      ConsumeToken();
+    } else if (TypeModifier == OMPC_MAP_MODIFIER_mapper) {
+      Data.MapTypeModifiers.push_back(TypeModifier);
+      Data.MapTypeModifiersLoc.push_back(Tok.getLocation());
+      ConsumeToken();
+      // Parse '('.
+      BalancedDelimiterTracker T(*this, tok::l_paren, tok::colon);
+      if (T.expectAndConsume(diag::err_expected_lparen_after,
+                             getOpenMPSimpleClauseTypeName(
+                                 OMPC_map, OMPC_MAP_MODIFIER_mapper))) {
+        SkipUntil(tok::colon, tok::annot_pragma_openmp_end, StopBeforeMatch);
+        return true;
+      }
+      // Parse mapper-identifier
+      if (getLangOpts().CPlusPlus)
+        ParseOptionalCXXScopeSpecifier(Data.MapperIdScopeSpec,
+                                       /*ObjectType=*/nullptr,
+                                       /*EnteringContext=*/false);
+      if (Tok.isNot(tok::identifier) && Tok.isNot(tok::kw_default)) {
+        Diag(Tok.getLocation(), diag::err_omp_mapper_illegal_identifier);
+        SkipUntil(tok::colon, tok::annot_pragma_openmp_end, StopBeforeMatch);
+        return true;
+      }
+      auto &DeclNames = Actions.getASTContext().DeclarationNames;
+      Data.MapperId = DeclarationNameInfo(
+          DeclNames.getIdentifier(Tok.getIdentifierInfo()), Tok.getLocation());
+      ConsumeToken();
+      // Parse ')'.
+      T.consumeClose();
     } else {
       // For the case of unknown map-type-modifier or a map-type.
       // Map-type is followed by a colon; the function returns when it
       // encounters a token followed by a colon.
       if (Tok.is(tok::comma)) {
-        P.Diag(Tok, diag::err_omp_map_type_modifier_missing);
-        P.ConsumeToken();
+        Diag(Tok, diag::err_omp_map_type_modifier_missing);
+        ConsumeToken();
         continue;
       }
       // Potential map-type token as it is followed by a colon.
       if (PP.LookAhead(0).is(tok::colon))
-        return;
-      P.Diag(Tok, diag::err_omp_unknown_map_type_modifier);
-      P.ConsumeToken();
+        return false;
+      Diag(Tok, diag::err_omp_unknown_map_type_modifier);
+      ConsumeToken();
     }
-    if (P.getCurToken().is(tok::comma))
-      P.ConsumeToken();
+    if (getCurToken().is(tok::comma))
+      ConsumeToken();
   }
+  return false;
 }
 
 /// Checks if the token is a valid map-type.
 static OpenMPMapClauseKind isMapType(Parser &P) {
   Token Tok = P.getCurToken();
   // The map-type token can be either an identifier or the C++ delete keyword.
   if (!Tok.isOneOf(tok::identifier, tok::kw_delete))
     return OMPC_MAP_unknown;

 
 /// Parses clauses with list.
 bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
                                 OpenMPClauseKind Kind,
                                 SmallVectorImpl<Expr *> &Vars,
                                 OpenMPVarListDataTy &Data) {
   UnqualifiedId UnqualifiedReductionId;
   bool InvalidReductionId = false;
+  bool IsInvalidMapModifier = false;
 
   // Parse '('.
   BalancedDelimiterTracker T(*this, tok::l_paren, tok::annot_pragma_openmp_end);
   if (T.expectAndConsume(diag::err_expected_lparen_after,
                          getOpenMPClauseName(Kind)))
     return true;
 
   bool NeedRParenForLinear = false;
if (SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end,
         StopBeforeMatch)) {
       if (Tok.is(tok::colon))
         ColonPresent = true;
     }
     TPA.Revert();
     // Only parse map-type-modifier[s] and map-type if a colon is present in
     // the map clause.
     if (ColonPresent) {
-      parseMapTypeModifiers(*this, Data);
+      IsInvalidMapModifier = parseMapTypeModifiers(Data);
+      if (!IsInvalidMapModifier)
         parseMapType(*this, Data);
     }
     if (Data.MapType == OMPC_MAP_unknown) {
       Data.MapType = OMPC_MAP_tofrom;
       Data.IsMapTypeImplicit = true;
     }
 
     if (Tok.is(tok::colon))
       Data.ColonLoc = ConsumeToken();
   }
bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
 
   // Parse ')'.
   Data.RLoc = Tok.getLocation();
   if (!T.consumeClose())
     Data.RLoc = T.getCloseLocation();
   return (Kind == OMPC_depend && Data.DepKind != OMPC_DEPEND_unknown &&
           Vars.empty()) ||
          (Kind != OMPC_depend && Kind != OMPC_map && Vars.empty()) ||
-         (MustHaveTail && !Data.TailExpr) || InvalidReductionId;
+         (MustHaveTail && !Data.TailExpr) || InvalidReductionId ||
+         IsInvalidMapModifier;
 }
 
 /// Parsing of OpenMP clause 'private', 'firstprivate', 'lastprivate',
 /// 'shared', 'copyin', 'copyprivate', 'flush', 'reduction', 'task_reduction' or
 /// 'in_reduction'.
 ///
 ///    private-clause:
 ///       'private' '(' list ')'

 ///    copyprivate-clause:
 ///       'copyprivate' '(' list ')'
 ///    flush-clause:
 ///       'flush' '(' list ')'
 ///    depend-clause:
 ///       'depend' '(' in | out | inout : list | source ')'
 ///    map-clause:
 ///       'map' '(' [ [ always [,] ] [ close [,] ]
+///          [ mapper(mapper-identifier) [,] ]
 ///          to | from | tofrom | alloc | release | delete ':' ] list ')';
 ///    to-clause:
 ///       'to' '(' list ')'
 ///    from-clause:
 ///       'from' '(' list ')'
 ///    use_device_ptr-clause:
 ///       'use_device_ptr' '(' list ')'
 ///    is_device_ptr-clause:
OMPClause *Parser::ParseOpenMPVarListClause(OpenMPDirectiveKind DKind,
   if (ParseOpenMPVarList(DKind, Kind, Vars, Data))
     return nullptr;
 
   if (ParseOnly)
     return nullptr;
   return Actions.ActOnOpenMPVarListClause(
       Kind, Vars, Data.TailExpr, Loc, LOpen, Data.ColonLoc, Data.RLoc,
       Data.ReductionIdScopeSpec, Data.ReductionId, Data.DepKind, Data.LinKind,
-      Data.MapTypeModifiers, Data.MapTypeModifiersLoc, Data.MapType,
-      Data.IsMapTypeImplicit, Data.DepLinMapLoc);
+      Data.MapTypeModifiers, Data.MapTypeModifiersLoc, Data.MapperIdScopeSpec,
+      Data.MapperId, Data.MapType, Data.IsMapTypeImplicit, Data.DepLinMapLoc);
 }
 

lib/Sema/SemaLookup.cpp

bool Sema::CppLookupName(LookupResult &R, Scope *S) {
   // Implicitly declare member functions with the name we're looking for, if in
   // fact we are in a scope where it matters.
 
   Scope *Initial = S;
   IdentifierResolver::iterator
     I = IdResolver.begin(Name),
     IEnd = IdResolver.end();
 
+  if (NameKind == LookupOMPMapperName) {
+    // Skip out-of-scope declarations.
+    for (; I != IEnd; ++I) {
+      Scope *CurS = S;
+      while (CurS && !CurS->isDeclScope(*I))
+        CurS = CurS->getParent();
+      if (CurS)
+        break;
+    }
+  }
+
   // First we lookup local scope.
   // We don't consider using-directives, as per 7.3.4.p1 [namespace.udir]
   // ...During unqualified name lookup (3.4.1), the names appear as if
   // they were declared in the nearest enclosing namespace which contains
   // both the using-directive and the nominated namespace.
   // [Note: in this context, "contains" means "contains directly or
   // indirectly".
   //
for (IdentifierResolver::iterator I = IdResolver.begin(Name),
           if (LeftStartingScope && !((*I)->hasLinkage())) {
             R.setShadowed();
             continue;
           }
         }
         else if (NameKind == LookupObjCImplicitSelfParam &&
                  !isa<ImplicitParamDecl>(*I))
           continue;
+        else if (NameKind == LookupOMPMapperName) {
+          // Skip out-of-scope declarations.
+          Scope *CurS = S;
+          while (CurS && !CurS->isDeclScope(*I))
+            CurS = CurS->getParent();
+          if (!CurS)
+            continue;
+        }
 
         R.addDecl(D);
 
         // Check whether there are any other declarations with the same name
         // and in the same scope.
         if (I != IEnd) {
           // Find the scope in which this declaration was declared (if it
           // actually exists in a Scope).

lib/Sema/SemaOpenMP.cpp

if (!ImplicitFirstprivates.empty()) {
         ClausesWithImplicit.push_back(Implicit);
         ErrorFound = cast<OMPFirstprivateClause>(Implicit)->varlist_size() !=
                      ImplicitFirstprivates.size();
       } else {
         ErrorFound = true;
       }
     }
     if (!ImplicitMaps.empty()) {
+      CXXScopeSpec MapperIdScopeSpec;
+      DeclarationNameInfo MapperId;
       if (OMPClause *Implicit = ActOnOpenMPMapClause(
-              llvm::None, llvm::None, OMPC_MAP_tofrom,
-              /*IsMapTypeImplicit=*/true, SourceLocation(), SourceLocation(),
-              ImplicitMaps, SourceLocation(), SourceLocation(),
-              SourceLocation())) {
+              llvm::None, llvm::None, MapperIdScopeSpec, MapperId,
+              OMPC_MAP_tofrom, /*IsMapTypeImplicit=*/true, SourceLocation(),
+              SourceLocation(), ImplicitMaps, SourceLocation(),
+              SourceLocation(), SourceLocation())) {
         ClausesWithImplicit.emplace_back(Implicit);
         ErrorFound |=
             cast<OMPMapClause>(Implicit)->varlist_size() != ImplicitMaps.size();
       } else {
         ErrorFound = true;
       }
     }
   }

 OMPClause *Sema::ActOnOpenMPVarListClause(
     OpenMPClauseKind Kind, ArrayRef<Expr *> VarList, Expr *TailExpr,
     SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation ColonLoc,
     SourceLocation EndLoc, CXXScopeSpec &ReductionIdScopeSpec,
     const DeclarationNameInfo &ReductionId, OpenMPDependClauseKind DepKind,
     OpenMPLinearClauseKind LinKind,
     ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
     ArrayRef<SourceLocation> MapTypeModifiersLoc,
+    CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId,
     OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
     SourceLocation DepLinMapLoc) {
   OMPClause *Res = nullptr;
   switch (Kind) {
   case OMPC_private:
     Res = ActOnOpenMPPrivateClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_firstprivate:
OMPClause *Sema::ActOnOpenMPVarListClause(
   case OMPC_flush:
     Res = ActOnOpenMPFlushClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_depend:
     Res = ActOnOpenMPDependClause(DepKind, DepLinMapLoc, ColonLoc, VarList,
                                   StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_map:
-    Res = ActOnOpenMPMapClause(MapTypeModifiers, MapTypeModifiersLoc, MapType,
+    Res = ActOnOpenMPMapClause(MapTypeModifiers, MapTypeModifiersLoc,
+                               MapperIdScopeSpec, MapperId, MapType,
                                IsMapTypeImplicit, DepLinMapLoc, ColonLoc,
                                VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_to:
     Res = ActOnOpenMPToClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_from:
     Res = ActOnOpenMPFromClause(VarList, StartLoc, LParenLoc, EndLoc);
ExprResult TransformMemberExpr(MemberExpr *E) {
     }
     return BaseTransform::TransformMemberExpr(E);
   }
   DeclRefExpr *getCapturedExpr() { return CapturedExpr; }
 };
 } // namespace
 
 template <typename T, typename U>
-static T filterLookupForUDR(SmallVectorImpl<U> &Lookups,
-                            const llvm::function_ref<T(ValueDecl *)> Gen) {
+static T filterLookupForUDReductionAndMapper(
+    SmallVectorImpl<U> &Lookups, const llvm::function_ref<T(ValueDecl *)> Gen) {
   for (U &Set : Lookups) {
     for (auto *D : Set) {
       if (T Res = Gen(cast<ValueDecl>(D)))
         return Res;
     }
   }
   return T();
 }
for (NamedDecl *D : ULE->decls()) {
       else if (auto *DRD = cast<OMPDeclareReductionDecl>(D))
         Lookups.back().addDecl(DRD);
       PrevD = D;
     }
   }
   if (SemaRef.CurContext->isDependentContext() || Ty->isDependentType() ||
       Ty->isInstantiationDependentType() ||
       Ty->containsUnexpandedParameterPack() ||
-      filterLookupForUDR<bool>(Lookups, [](ValueDecl *D) {
+      filterLookupForUDReductionAndMapper<bool>(Lookups, [](ValueDecl *D) {
         return !D->isInvalidDecl() &&
                (D->getType()->isDependentType() ||
                 D->getType()->isInstantiationDependentType() ||
                 D->getType()->containsUnexpandedParameterPack());
       })) {
     UnresolvedSet<8> ResSet;
     for (const UnresolvedSet<8> &Set : Lookups) {
       if (Set.empty())
if (SemaRef.isCompleteType(Loc, Ty) || TyRec->isBeingDefined() ||
       if (Lookup.empty()) {
         Lookups.emplace_back();
         Lookups.back().append(Lookup.begin(), Lookup.end());
       }
     }
   }
   // Perform ADL.
   argumentDependentLookup(SemaRef, ReductionId, Loc, Ty, Lookups);
-  if (auto *VD = filterLookupForUDR<ValueDecl *>(
+  if (auto *VD = filterLookupForUDReductionAndMapper<ValueDecl *>(
           Lookups, [&SemaRef, Ty](ValueDecl *D) -> ValueDecl * {
             if (!D->isInvalidDecl() &&
                 SemaRef.Context.hasSameType(D->getType(), Ty))
               return D;
             return nullptr;
           }))
     return SemaRef.BuildDeclRefExpr(VD, VD->getType().getNonReferenceType(),
                                     VK_LValue, Loc);
-  if (auto *VD = filterLookupForUDR<ValueDecl *>(
+  if (auto *VD = filterLookupForUDReductionAndMapper<ValueDecl *>(
           Lookups, [&SemaRef, Ty, Loc](ValueDecl *D) -> ValueDecl * {
             if (!D->isInvalidDecl() &&
                 SemaRef.IsDerivedFrom(Loc, Ty, D->getType()) &&
                 !Ty.isMoreQualifiedThan(D->getType()))
               return D;
             return nullptr;
           })) {
     CXXBasePaths Paths(/*FindAmbiguities=*/true, /*RecordPaths=*/true,
if (EnclosingExpr && !IsEnclosedByDataEnvironmentExpr) {
     SemaRef.Diag(EnclosingExpr->getExprLoc(), diag::note_used_here)
         << EnclosingExpr->getSourceRange();
     return true;
   }
 
   return FoundError;
 }
 
+// Look up the user-defined mapper given the mapper name and mapped type, and
+// build a reference to it.
+ExprResult buildUserDefinedMapperRef(Sema &SemaRef, Scope *S,
+                                     CXXScopeSpec &MapperIdScopeSpec,
+                                     DeclarationNameInfo &MapperId,
+                                     QualType Type, Expr *UnresolvedMapper) {
+  if (MapperIdScopeSpec.isInvalid())
+    return ExprError();
+  // Find all user-defined mappers with the given MapperId.
+  SmallVector<UnresolvedSet<8>, 4> Lookups;
+  LookupResult Lookup(SemaRef, MapperId, Sema::LookupOMPMapperName);
+  Lookup.suppressDiagnostics();
+  if (S) {
+    while (S && SemaRef.LookupParsedName(Lookup, S, &MapperIdScopeSpec)) {
+      NamedDecl *D = Lookup.getRepresentativeDecl();
+      while (S && !S->isDeclScope(D))
+        S = S->getParent();
+      if (S)
+        S = S->getParent();
+      Lookups.emplace_back();
+      Lookups.back().append(Lookup.begin(), Lookup.end());
+      Lookup.clear();
+    }
+  } else if (auto *ULE = cast_or_null<UnresolvedLookupExpr>(UnresolvedMapper)) {
+    // Extract the user-defined mappers with the given MapperId.
+    Lookups.push_back(UnresolvedSet<8>());
+    for (NamedDecl *D : ULE->decls()) {
+      auto *DMD = cast<OMPDeclareMapperDecl>(D);
+      assert(DMD && "Expect valid OMPDeclareMapperDecl during instantiation.");
+      Lookups.back().addDecl(DMD);
+    }
+  }
+  // Defer the lookup for dependent types. The results will be passed through
+  // UnresolvedMapper on instantiation.
+  if (SemaRef.CurContext->isDependentContext() || Type->isDependentType() ||
+      Type->isInstantiationDependentType() ||
+      Type->containsUnexpandedParameterPack() ||
+      filterLookupForUDReductionAndMapper<bool>(Lookups, [](ValueDecl *D) {
+        return !D->isInvalidDecl() &&
+               (D->getType()->isDependentType() ||
+                D->getType()->isInstantiationDependentType() ||
+                D->getType()->containsUnexpandedParameterPack());
+      })) {
+    UnresolvedSet<8> URS;
+    for (const UnresolvedSet<8> &Set : Lookups) {
+      if (Set.empty())
+        continue;
+      URS.append(Set.begin(), Set.end());
+    }
+    return UnresolvedLookupExpr::Create(
+        SemaRef.Context, /*NamingClass=*/nullptr,
+        MapperIdScopeSpec.getWithLocInContext(SemaRef.Context), MapperId,
+        /*ADL=*/false, /*Overloaded=*/true, URS.begin(), URS.end());
+  }
+  // [OpenMP 5.0], 2.19.7.3 declare mapper Directive, Restrictions
+  //  The type must be of struct, union or class type in C and C++
+  if (!Type->isStructureOrClassType() && !Type->isUnionType())
+    return ExprEmpty();
+  // Return the first user-defined mapper with the desired type.
+  SourceLocation Loc = MapperId.getLoc();
+  if (auto *VD = filterLookupForUDReductionAndMapper<ValueDecl *>(
+          Lookups, [&SemaRef, Type](ValueDecl *D) -> ValueDecl * {
+            if (!D->isInvalidDecl() &&
+                SemaRef.Context.hasSameType(D->getType(), Type))
+              return D;
+            return nullptr;
+          }))
+    return SemaRef.BuildDeclRefExpr(VD, Type, VK_LValue, Loc);
+  // Find the first user-defined mapper with a type derived from the desired
+  // type.
+  if (auto *VD = filterLookupForUDReductionAndMapper<ValueDecl *>(
+          Lookups, [&SemaRef, Type, Loc](ValueDecl *D) -> ValueDecl * {
+            if (!D->isInvalidDecl() &&
+                SemaRef.IsDerivedFrom(Loc, Type, D->getType()) &&
+                !Type.isMoreQualifiedThan(D->getType()))
+              return D;
+            return nullptr;
+          })) {
+    CXXBasePaths Paths(/*FindAmbiguities=*/true, /*RecordPaths=*/true,
+                       /*DetectVirtual=*/false);
+    if (SemaRef.IsDerivedFrom(Loc, Type, VD->getType(), Paths)) {
+      if (!Paths.isAmbiguous(SemaRef.Context.getCanonicalType(
+              VD->getType().getUnqualifiedType()))) {
+        if (SemaRef.CheckBaseClassAccess(
+                Loc, VD->getType(), Type, Paths.front(),
+                /*DiagID=*/0) != Sema::AR_inaccessible) {
+          return SemaRef.BuildDeclRefExpr(VD, Type, VK_LValue, Loc);
+        }
+      }
+    }
+  }
+  // Report error if a mapper is specified, but cannot be found.
+  if (MapperIdScopeSpec.isSet() || MapperId.getAsString() != "default") {
+    SemaRef.Diag(Loc, diag::err_omp_invalid_mapper)
+        << Type << MapperId.getName();
+    return ExprError();
+  }
+  return ExprEmpty();
+}
+
 namespace {
 // Utility struct that gathers all the related lists associated with a mappable
 // expression.
 struct MappableVarListInfo {
   // The list of expressions.
   ArrayRef<Expr *> VarList;
   // The list of processed expressions.
   SmallVector<Expr *, 16> ProcessedVarList;
   // The mappble components for each expression.
   OMPClauseMappableExprCommon::MappableExprComponentLists VarComponents;
   // The base declaration of the variable.
   SmallVector<ValueDecl *, 16> VarBaseDeclarations;
+  // The reference to the user-defined mapper associated with every expression.
+  SmallVector<Expr *, 16> UDMapperList;
 
   MappableVarListInfo(ArrayRef<Expr *> VarList) : VarList(VarList) {
     // We have a list of components and base declarations for each entry in the
     // variable list.
     VarComponents.reserve(VarList.size());
     VarBaseDeclarations.reserve(VarList.size());
   }
 };
 }
 
 // Check the validity of the provided variable list for the provided clause kind
-// \a CKind. In the check process the valid expressions, and mappable expression
-// components and variables are extracted and used to fill \a Vars,
-// \a ClauseComponents, and \a ClauseBaseDeclarations. \a MapType and
-// \a IsMapTypeImplicit are expected to be valid if the clause kind is 'map'.
-static void
-checkMappableExpressionList(Sema &SemaRef, DSAStackTy *DSAS,
-                            OpenMPClauseKind CKind, MappableVarListInfo &MVLI,
-                            SourceLocation StartLoc,
-                            OpenMPMapClauseKind MapType = OMPC_MAP_unknown,
-                            bool IsMapTypeImplicit = false) {
+// \a CKind. In the check process the valid expressions, mappable expression
+// components, variables, and user-defined mappers are extracted and used to
+// fill \a ProcessedVarList, \a VarComponents, \a VarBaseDeclarations, and \a
+// UDMapperList in MVLI. \a MapType, \a IsMapTypeImplicit, \a MapperIdScopeSpec,
+// and \a MapperId are expected to be valid if the clause kind is 'map'.
+static void checkMappableExpressionList(
+    Sema &SemaRef, DSAStackTy *DSAS, OpenMPClauseKind CKind,
+    MappableVarListInfo &MVLI, SourceLocation StartLoc,
+    bool IsMapTypeImplicit = false, CXXScopeSpec *MapperIdScopeSpec = nullptr,
+    DeclarationNameInfo *MapperId = nullptr,
+    ArrayRef<Expr *> UnresolvedMappers = llvm::None) {
   // We only expect mappable expressions in 'to', 'from', and 'map' clauses.
   assert((CKind == OMPC_map || CKind == OMPC_to || CKind == OMPC_from) &&
          "Unexpected clause kind with mappable expressions!");
+  assert(
+      ((CKind == OMPC_map && MapperIdScopeSpec && MapperId) ||
+       (CKind != OMPC_map && !MapperIdScopeSpec && !MapperId)) &&
+      "Map clauses and only map clauses have user-defined mapper identifiers.");
+
+  // Iterators to find the current unresolved mapper expression.
+  auto UMIt = UnresolvedMappers.begin(), UMEnd = UnresolvedMappers.end();
+  bool UpdateUMIt = false;
+  Expr *UnresolvedMapper = nullptr;
 
   // Keep track of the mappable components and base declarations in this clause.
   // Each entry in the list is going to have a list of components associated. We
   // record each set of the components so that we can build the clause later on.
   // In the end we should have the same amount of declarations and component
   // lists.
 
   for (Expr *RE : MVLI.VarList) {
     assert(RE && "Null expr in omp to/from/map clause");
     SourceLocation ELoc = RE->getExprLoc();
 
+    // Find the current unresolved mapper expression.
+    if (UpdateUMIt && UMIt != UMEnd) {
+      UMIt++;
+      std::cerr << "Not empty.\n";
+      assert(
+          UMIt != UMEnd &&
+          "Expect the size of UnresolvedMappers to match with that of VarList");
+    }
+    UpdateUMIt = true;
+    if (UMIt != UMEnd)
+      UnresolvedMapper = *UMIt;
+
     const Expr *VE = RE->IgnoreParenLValueCasts();
 
     if (VE->isValueDependent() || VE->isTypeDependent() ||
         VE->isInstantiationDependent() ||
         VE->containsUnexpandedParameterPack()) {
+      if (CKind == OMPC_map) {
+        // Try to find the associated user-defined mapper.
+        ExprResult ER = buildUserDefinedMapperRef(
+            SemaRef, DSAS->getCurScope(), *MapperIdScopeSpec, *MapperId,
+            VE->getType().getCanonicalType(), UnresolvedMapper);
+        if (ER.isInvalid()) {
+          continue;
+        } else if (ER.isUsable()) {
+          MVLI.UDMapperList.push_back(ER.get());
+        } else {
+          MVLI.UDMapperList.push_back(nullptr);
+        }
+      } else {
+        MVLI.UDMapperList.push_back(nullptr);
+      }
       // 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_map) {
+        // Try to find the associated user-defined mapper.
+        ExprResult ER = buildUserDefinedMapperRef(
+            SemaRef, DSAS->getCurScope(), *MapperIdScopeSpec, *MapperId,
+            VE->getType().getCanonicalType(), UnresolvedMapper);
+        if (ER.isInvalid()) {
+          continue;
+        } else if (ER.isUsable()) {
+          MVLI.UDMapperList.push_back(ER.get());
+        } else {
+          MVLI.UDMapperList.push_back(nullptr);
+        }
+      } else {
+        MVLI.UDMapperList.push_back(nullptr);
+      }
       // Skip restriction checking for variable or field declarations
       MVLI.ProcessedVarList.push_back(RE);
       MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
       MVLI.VarComponents.back().append(CurComponents.begin(),
                                        CurComponents.end());
       MVLI.VarBaseDeclarations.push_back(nullptr);
       continue;
     }
if (CKind == OMPC_map) {
           SemaRef.Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa)
               << getOpenMPClauseName(DVar.CKind)
               << getOpenMPClauseName(OMPC_map)
               << getOpenMPDirectiveName(DSAS->getCurrentDirective());
           reportOriginalDsa(SemaRef, DSAS, CurDeclaration, DVar);
           continue;
         }
       }
+
+      // Try to find the associated user-defined mapper.
+      ExprResult ER = buildUserDefinedMapperRef(
+          SemaRef, DSAS->getCurScope(), *MapperIdScopeSpec, *MapperId,
+          Type.getCanonicalType(), UnresolvedMapper);
+      if (ER.isInvalid()) {
+        continue;
+      } else if (ER.isUsable()) {
+        MVLI.UDMapperList.push_back(ER.get());
+      } else {
+        MVLI.UDMapperList.push_back(nullptr);
+      }
+    } else {
+      MVLI.UDMapperList.push_back(nullptr);
     }
 
     // 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);
 
     // Save the components and declaration to create the clause. For purposes of
     // the clause creation, any component list that has has base 'this' uses
     // null as base declaration.
     MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
     MVLI.VarComponents.back().append(CurComponents.begin(),
                                      CurComponents.end());
     MVLI.VarBaseDeclarations.push_back(isa<MemberExpr>(BE) ? nullptr
                                                            : CurDeclaration);
   }
 }
 
-OMPClause *
-Sema::ActOnOpenMPMapClause(ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
-                           ArrayRef<SourceLocation> MapTypeModifiersLoc,
-                           OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
-                           SourceLocation MapLoc, SourceLocation ColonLoc,
-                           ArrayRef<Expr *> VarList, SourceLocation StartLoc,
-                           SourceLocation LParenLoc, SourceLocation EndLoc) {
-  MappableVarListInfo MVLI(VarList);
-  checkMappableExpressionList(*this, DSAStack, OMPC_map, MVLI, StartLoc,
-                              MapType, IsMapTypeImplicit);
-
-  OpenMPMapModifierKind Modifiers[] = { OMPC_MAP_MODIFIER_unknown,
-                                        OMPC_MAP_MODIFIER_unknown };
+OMPClause *Sema::ActOnOpenMPMapClause(
+    ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
+    CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId,
+    OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, SourceLocation MapLoc,
+    SourceLocation ColonLoc, ArrayRef<Expr *> VarList, SourceLocation StartLoc,
+    SourceLocation LParenLoc, SourceLocation EndLoc,
+    ArrayRef<Expr *> UnresolvedMappers) {
+  OpenMPMapModifierKind Modifiers[] = {OMPC_MAP_MODIFIER_unknown,
+                                       OMPC_MAP_MODIFIER_unknown,
+                                       OMPC_MAP_MODIFIER_unknown};
   SourceLocation ModifiersLoc[OMPMapClause::NumberOfModifiers];
 
   // Process map-type-modifiers, flag errors for duplicate modifiers.
   unsigned Count = 0;
   for (unsigned I = 0, E = MapTypeModifiers.size(); I < E; ++I) {
     if (MapTypeModifiers[I] != OMPC_MAP_MODIFIER_unknown &&
         llvm::find(Modifiers, MapTypeModifiers[I]) != std::end(Modifiers)) {
       Diag(MapTypeModifiersLoc[I], diag::err_omp_duplicate_map_type_modifier);
       continue;
     }
     assert(Count < OMPMapClause::NumberOfModifiers &&
            "Modifiers exceed the allowed number of map type modifiers");
     Modifiers[Count] = MapTypeModifiers[I];
     ModifiersLoc[Count] = MapTypeModifiersLoc[I];
     ++Count;
   }
 
+  // If the identifier of user-defined mapper is not specified, it is "default".
+  if (MapperId.getName().isEmpty()) {
+    auto &DeclNames = getASTContext().DeclarationNames;
+    MapperId.setName(
+        DeclNames.getIdentifier(&getASTContext().Idents.get("default")));
+  }
+
+  MappableVarListInfo MVLI(VarList);
+  checkMappableExpressionList(*this, DSAStack, OMPC_map, MVLI, StartLoc,
+                              MapType, IsMapTypeImplicit, &MapperIdScopeSpec,
+                              &MapperId, UnresolvedMappers);
+
   // We need to produce a map clause even if we don't have variables so that
   // other diagnostics related with non-existing map clauses are accurate.
-  return OMPMapClause::Create(Context, StartLoc, LParenLoc, EndLoc,
-                              MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
-                              MVLI.VarComponents, Modifiers, ModifiersLoc,
-                              MapType, IsMapTypeImplicit, MapLoc);
+  return OMPMapClause::Create(
+      Context, StartLoc, LParenLoc, EndLoc, MVLI.ProcessedVarList,
+      MVLI.VarBaseDeclarations, MVLI.VarComponents, MVLI.UDMapperList,
+      Modifiers, ModifiersLoc, MapperIdScopeSpec.getWithLocInContext(Context),
+      MapperId, MapType, IsMapTypeImplicit, MapLoc);
 }
 
 QualType Sema::ActOnOpenMPDeclareReductionType(SourceLocation TyLoc,
                                                TypeResult ParsedType) {
   assert(ParsedType.isUsable());
 
   QualType ReductionType = GetTypeFromParser(ParsedType.get());
   if (ReductionType.isNull())

lib/Sema/SemaTemplateInstantiateDecl.cpp

for (OMPClause *C : D->clauselists()) {
         if (!NE) {
           IsCorrect = false;
           break;
         }
         NewVars.push_back(NE);
       }
       if (!IsCorrect)
         break;
+      NestedNameSpecifierLoc NewQualifierLoc =
+          SemaRef.SubstNestedNameSpecifierLoc(OldC->getMapperQualifierLoc(),
+                                              TemplateArgs);
+      CXXScopeSpec SS;
+      SS.Adopt(NewQualifierLoc);
+      DeclarationNameInfo NewNameInfo = SemaRef.SubstDeclarationNameInfo(
+          OldC->getMapperIdInfo(), TemplateArgs);
       OMPClause *NewC = SemaRef.ActOnOpenMPMapClause(
           OldC->getMapTypeModifiers(), OldC->getMapTypeModifiersLoc(),
-          OldC->getMapType(), OldC->isImplicitMapType(), OldC->getMapLoc(),
-          OldC->getColonLoc(), NewVars, OldC->getBeginLoc(),
-          OldC->getLParenLoc(), OldC->getEndLoc());
+          SS, NewNameInfo, OldC->getMapType(),
+          OldC->isImplicitMapType(), OldC->getMapLoc(), OldC->getColonLoc(),
+          NewVars, OldC->getBeginLoc(), OldC->getLParenLoc(),
+          OldC->getEndLoc());
       Clauses.push_back(NewC);
     }
     SemaRef.EndOpenMPDSABlock(nullptr);
   }
   (void)SemaRef.ActOnOpenMPDeclareMapperDirectiveEnd(NewDMD, /*S=*/nullptr,
                                                      Clauses);
   if (!IsCorrect)
     return nullptr;

lib/Sema/TreeTransform.h

OMPClause *RebuildOMPDeviceClause(Expr *Device, SourceLocation StartLoc,
     return getSema().ActOnOpenMPDeviceClause(Device, StartLoc, LParenLoc,
                                              EndLoc);
   }
 
   /// Build a new OpenMP 'map' clause.
   ///
   /// By default, performs semantic analysis to build the new OpenMP clause.
   /// Subclasses may override this routine to provide different behavior.
-  OMPClause *
-  RebuildOMPMapClause(ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
-                      ArrayRef<SourceLocation> MapTypeModifiersLoc,
-                      OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
-                      SourceLocation MapLoc, SourceLocation ColonLoc,
-                      ArrayRef<Expr *> VarList, SourceLocation StartLoc,
-                      SourceLocation LParenLoc, SourceLocation EndLoc) {
-    return getSema().ActOnOpenMPMapClause(MapTypeModifiers, MapTypeModifiersLoc,
-                                          MapType, IsMapTypeImplicit, MapLoc,
-                                          ColonLoc, VarList, StartLoc,
-                                          LParenLoc, EndLoc);
+  OMPClause *RebuildOMPMapClause(
+      ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
+      CXXScopeSpec MapperIdScopeSpec, DeclarationNameInfo MapperId,
+      SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
+      SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc,
+      ArrayRef<Expr *> UnresolvedMappers) {
+    return getSema().ActOnOpenMPMapClause(
+        MapTypeModifiers, MapTypeModifiersLoc, MapperIdScopeSpec, MapperId,
+        MapType, IsMapTypeImplicit, MapLoc, ColonLoc, VarList, StartLoc,
+        LParenLoc, EndLoc, UnresolvedMappers);
   }
 
   /// Build a new OpenMP 'num_teams' clause.
   ///
   /// By default, performs semantic analysis to build the new statement.
   /// Subclasses may override this routine to provide different behavior.
   OMPClause *RebuildOMPNumTeamsClause(Expr *NumTeams, SourceLocation StartLoc,
                                       SourceLocation LParenLoc,
OMPClause *TreeTransform<Derived>::TransformOMPMapClause(OMPMapClause *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 nullptr;
     Vars.push_back(EVar.get());
   }
+  NestedNameSpecifierLoc QualifierLoc;
+  if (C->getMapperQualifierLoc()) {
+    QualifierLoc = getDerived().TransformNestedNameSpecifierLoc(
+        C->getMapperQualifierLoc());
+    if (!QualifierLoc)
+      return nullptr;
+  }
+  CXXScopeSpec MapperIdScopeSpec;
+  MapperIdScopeSpec.Adopt(QualifierLoc);
+  DeclarationNameInfo MapperIdInfo = C->getMapperIdInfo();
+  if (MapperIdInfo.getName()) {
+    MapperIdInfo = getDerived().TransformDeclarationNameInfo(MapperIdInfo);
+    if (!MapperIdInfo.getName())
+      return nullptr;
+  }
+  // Build a list of all candidate OMPDeclareMapperDecls, which is provided by
+  // the previous user-defined mapper lookup in dependent environment.
+  llvm::SmallVector<Expr *, 16> UnresolvedMappers;
+  for (auto *E : C->mapperlists()) {
+    // Transform all the decls.
+    if (E) {
+      auto *ULE = cast<UnresolvedLookupExpr>(E);
+      UnresolvedSet<8> Decls;
+      for (auto *D : ULE->decls()) {
+        NamedDecl *InstD =
+            cast<NamedDecl>(getDerived().TransformDecl(E->getExprLoc(), D));
+        Decls.addDecl(InstD, InstD->getAccess());
+      }
+      UnresolvedMappers.push_back(UnresolvedLookupExpr::Create(
+          SemaRef.Context, /*NamingClass=*/nullptr,
+          MapperIdScopeSpec.getWithLocInContext(SemaRef.Context), MapperIdInfo,
+          /*ADL=*/false, ULE->isOverloaded(), Decls.begin(), Decls.end()));
+    } else
+      UnresolvedMappers.push_back(nullptr);
+  }
   return getDerived().RebuildOMPMapClause(
-      C->getMapTypeModifiers(), C->getMapTypeModifiersLoc(), C->getMapType(),
-      C->isImplicitMapType(), C->getMapLoc(), C->getColonLoc(), Vars,
-      C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
+      C->getMapTypeModifiers(), C->getMapTypeModifiersLoc(), MapperIdScopeSpec,
+      MapperIdInfo, C->getMapType(), C->isImplicitMapType(), C->getMapLoc(),
+      C->getColonLoc(), Vars, C->getBeginLoc(), C->getLParenLoc(),
+      C->getEndLoc(), UnresolvedMappers);
 }
 
 template <typename Derived>
 OMPClause *
 TreeTransform<Derived>::TransformOMPNumTeamsClause(OMPNumTeamsClause *C) {
   ExprResult E = getDerived().TransformExpr(C->getNumTeams());
   if (E.isInvalid())
     return nullptr;

lib/Serialization/ASTReader.cpp

 
 void OMPClauseReader::VisitOMPMapClause(OMPMapClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   for (unsigned I = 0; I < OMPMapClause::NumberOfModifiers; ++I) {
     C->setMapTypeModifier(
         I, static_cast<OpenMPMapModifierKind>(Record.readInt()));
     C->setMapTypeModifierLoc(I, Record.readSourceLocation());
   }
+  C->setMapperQualifierLoc(Record.readNestedNameSpecifierLoc());
+  DeclarationNameInfo DNI;
+  Record.readDeclarationNameInfo(DNI);
+  C->setMapperIdInfo(DNI);
   C->setMapType(
      static_cast<OpenMPMapClauseKind>(Record.readInt()));
   C->setMapLoc(Record.readSourceLocation());
   C->setColonLoc(Record.readSourceLocation());
   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.readExpr());
   C->setVarRefs(Vars);
 
+  SmallVector<Expr *, 16> UDMappers;
+  UDMappers.reserve(NumVars);
+  for (unsigned i = 0; i < NumVars; ++i)
+    UDMappers.push_back(Record.readExpr());
+  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);

lib/Serialization/ASTWriter.cpp

void OMPClauseWriter::VisitOMPMapClause(OMPMapClause *C) {
   Record.push_back(C->getUniqueDeclarationsNum());
   Record.push_back(C->getTotalComponentListNum());
   Record.push_back(C->getTotalComponentsNum());
   Record.AddSourceLocation(C->getLParenLoc());
   for (unsigned I = 0; I < OMPMapClause::NumberOfModifiers; ++I) {
     Record.push_back(C->getMapTypeModifier(I));
     Record.AddSourceLocation(C->getMapTypeModifierLoc(I));
   }
+  Record.AddNestedNameSpecifierLoc(C->getMapperQualifierLoc());
+  Record.AddDeclarationNameInfo(C->getMapperIdInfo());
   Record.push_back(C->getMapType());
   Record.AddSourceLocation(C->getMapLoc());
   Record.AddSourceLocation(C->getColonLoc());
   for (auto *E : C->varlists())
     Record.AddStmt(E);
+  for (auto *E : C->mapperlists())
+    Record.AddStmt(E);
   for (auto *D : C->all_decls())
     Record.AddDeclRef(D);
   for (auto N : C->all_num_lists())
     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());

test/OpenMP/declare_mapper_ast_print.c

 
 // CHECK: int main() {
 int main() {
 #pragma omp declare mapper(id: struct vec v) map(v.len)
 // CHECK: #pragma omp declare mapper (id : struct vec v) map(tofrom: v.len)
   {
 #pragma omp declare mapper(id: struct vec v) map(v.len)
 // CHECK: #pragma omp declare mapper (id : struct vec v) map(tofrom: v.len)
+    struct vec vv;
+    struct dat dd[10];
+#pragma omp target map(mapper(id) alloc: vv)
+// CHECK: #pragma omp target map(mapper(id),alloc: vv)
+    { vv.len++; }
+#pragma omp target map(mapper(default) from: dd[0:10])
+// CHECK: #pragma omp target map(mapper(default),from: dd[0:10])
+    { dd[0].i++; }
   }
   return 0;
 }
 // CHECK: }
 
 #endif

test/OpenMP/declare_mapper_ast_print.cpp

 // CHECK: class vec {
 class vec {
 public:
   int len;
   double *data;
 };
 // CHECK: };
 
+// CHECK: class vecchild : public N1::vec {
+class vecchild : public vec {
+public:
+  int lenc;
+};
+// CHECK: };
+
 #pragma omp declare mapper(id: vec v) map(v.len)
 // CHECK: #pragma omp declare mapper (id : N1::vec v) map(tofrom: v.len){{$}}
 };
 // CHECK: }
 // CHECK: ;
 
 template <class T>
 class dat {

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

test/OpenMP/declare_mapper_messages.c

 
 int fun(int arg) {
 #pragma omp declare mapper(id: struct vec v) map(v.len)
   {
 #pragma omp declare mapper(id: struct vec v) map(v.len)                 // expected-note {{previous definition is here}}
 #pragma omp declare mapper(id: struct vec v) map(v.len)                 // expected-error {{redefinition of user-defined mapper for type 'struct vec' with name 'id'}}
     {
 #pragma omp declare mapper(id: struct vec v) map(v.len)
+      struct vec vv, v1;
+#pragma omp target map(mapper)                                          // expected-error {{use of undeclared identifier 'mapper'}}
+      {}
+#pragma omp target map(mapper:vv)                                       // expected-error {{expected '(' after 'mapper'}}
+      {}
+#pragma omp target map(mapper( :vv)                                     // expected-error {{expected expression}} expected-error {{expected ')'}} expected-warning {{implicit declaration of function 'mapper' is invalid in C99}} expected-note {{to match this '('}}
+      {}
+#pragma omp target map(mapper(aa :vv)                                   // expected-error {{use of undeclared identifier 'aa'}} expected-error {{expected ')'}} expected-warning {{implicit declaration of function 'mapper' is invalid in C99}} expected-note {{to match this '('}}
+      {}
+#pragma omp target map(mapper(ab) :vv)                                  // expected-error {{missing map type}} expected-error {{cannot find a valid user-defined mapper for type 'struct vec' with name 'ab'}}
+      {}
+#pragma omp target map(mapper(aa) :vv)                                  // expected-error {{missing map type}}
+      {}
+#pragma omp target map(mapper(aa) to:vv) map(close mapper(aa) from:v1)
+      {}
     }
   }
   return arg;
 }

test/OpenMP/declare_mapper_messages.cpp

 }
 
 int fun(int arg) {
 #pragma omp declare mapper(id: vec v) map(v.len)
   {
 #pragma omp declare mapper(id: vec v) map(v.len)                        // expected-note {{previous definition is here}}
     {
 #pragma omp declare mapper(id: vec v) map(v.len)
+      vec vv, v1;
+#pragma omp target map(mapper)                                          // expected-error {{use of undeclared identifier 'mapper'}}
+      {}
+#pragma omp target map(mapper:vv)                                       // expected-error {{expected '(' after 'mapper'}}
+      {}
+#pragma omp target map(mapper( :vv)                                     // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+      {}
+#pragma omp target map(mapper(aa :vv)                                   // expected-error {{use of undeclared identifier 'aa'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+      {}
+#pragma omp target map(mapper(ab) :vv)                                  // expected-error {{missing map type}} expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'ab'}}
+      {}
+#pragma omp target map(mapper(N2::) :vv)                                // expected-error {{use of undeclared identifier 'N2'}} expected-error {{illegal identifier on 'omp declare mapper' directive}}
+      {}
+#pragma omp target map(mapper(N1::) :vv)                                // expected-error {{illegal identifier on 'omp declare mapper' directive}}
+      {}
+#pragma omp target map(mapper(aa) :vv)                                  // expected-error {{missing map type}}
+      {}
+#pragma omp target map(mapper(N1::aa) alloc:vv)                         // expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'aa'}}
+      {}
+#pragma omp target map(mapper(aa) to:vv) map(close mapper(aa) from:v1)
+      {}
+#pragma omp target map(mapper(N1::stack<int>::id) to:vv)
+      {}
     }
 #pragma omp declare mapper(id: vec v) map(v.len)                        // expected-error {{redefinition of user-defined mapper for type 'vec' with name 'id'}}
   }
   return arg;
 }

test/OpenMP/target_map_messages.cpp

void func(int arg) {
     #pragma omp target map(always, close, always, close, tofrom: a)   // expected-error {{same map type modifier has been specified more than once}} expected-error {{same map type modifier has been specified more than once}}
     {}
     #pragma omp target map( , tofrom: a)   // expected-error {{missing map type modifier}}
     {}
     #pragma omp target map( , , tofrom: a)   // expected-error {{missing map type modifier}} expected-error {{missing map type modifier}}
     {}
     #pragma omp target map( , , : a)   // expected-error {{missing map type modifier}} expected-error {{missing map type modifier}} expected-error {{missing map type}}
     {}
-    #pragma omp target map( d, f, bf: a)   // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+    #pragma omp target map( d, f, bf: a)   // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
     {}
-    #pragma omp target map( , f, : a)   // expected-error {{missing map type modifier}} expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+    #pragma omp target map( , f, : a)   // expected-error {{missing map type modifier}} expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
     {}
     #pragma omp target map(always close: a)   // expected-error {{missing map type}}
     {}
     #pragma omp target map(always close bf: a)   // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
     {}
-    #pragma omp target map(always tofrom close: a)   // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+    #pragma omp target map(always tofrom close: a)   // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
     {}
-    #pragma omp target map(tofrom from: a)   // expected-error {{incorrect map type modifier, expected 'always' or 'close'}}
+    #pragma omp target map(tofrom from: a)   // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
     {}
     #pragma omp target map(close bf: a)   // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
     {}
     return;
   }
 };
 
 struct SB {

 #pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
 #pragma omp target data map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
 #pragma omp target map(l)
   foo();
 
 #pragma omp target data map(always, tofrom: x)
 #pragma omp target data map(always: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 
 #pragma omp target data map(close, tofrom: x)
 #pragma omp target data map(close: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
 #pragma omp target data map(close, tofrom: close, tofrom, x)
   foo();
   return 0;
 }
 
 struct SA1{
   int a;
   struct SA1 *p;

 #pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
 #pragma omp target data map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
 #pragma omp target map(l)
   foo();
 
 #pragma omp target data map(always, tofrom: x)
 #pragma omp target data map(always: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 #pragma omp target data map(close, tofrom: x)
 #pragma omp target data map(close: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   foo();
 #pragma omp target private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target' directive}}  expected-note {{defined as private}}
   {}
 #pragma omp target firstprivate(j) map(j)  // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target' directive}} expected-note {{defined as firstprivate}}
   {}
 #pragma omp target map(m)
   {}
 // expected-note@+1 {{used here}}

test/OpenMP/target_parallel_for_map_messages.cpp

 #pragma omp target parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
 }
 

test/OpenMP/target_parallel_for_simd_map_messages.cpp

 #pragma omp target parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
 }
 

test/OpenMP/target_parallel_map_messages.cpp

 #pragma omp target parallel map(l)
   foo();
 }
 
 #pragma omp target parallel map(always, tofrom: x)
   foo();
 #pragma omp target parallel map(always: x) // expected-error {{missing map type}}
   foo();
-#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   foo();
 #pragma omp target parallel map(always, tofrom: always, tofrom, x)
   foo();
 #pragma omp target parallel map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 
   return 0;
 }

 #pragma omp target parallel map(l)
   foo();
 }
 
 #pragma omp target parallel map(always, tofrom: x)
   foo();
 #pragma omp target parallel map(always: x) // expected-error {{missing map type}}
   foo();
-#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   foo();
 #pragma omp target parallel map(always, tofrom: always, tofrom, x)
   foo();
 #pragma omp target parallel map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 
   return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
 }
 

test/OpenMP/target_simd_map_messages.cpp

 #pragma omp target simd map(l) // OK
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target simd map(l) //
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
 }
 

test/OpenMP/target_teams_distribute_map_messages.cpp

 #pragma omp target teams distribute map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target teams distribute map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
 }
 

test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp

 #pragma omp target teams distribute parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target teams distribute parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
 }
 

test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp

 #pragma omp target teams distribute parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target teams distribute parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
 }
 

test/OpenMP/target_teams_distribute_simd_map_messages.cpp

 #pragma omp target teams distribute simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target teams distribute simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
 }
 

test/OpenMP/target_teams_map_messages.cpp

 #pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
 #pragma omp target data map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
 #pragma omp target teams map(l)
   foo();
 
 #pragma omp target data map(always, tofrom: x)
 #pragma omp target data map(always: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
   return 0;
 }
 
 int main(int argc, char **argv) {
   const int d = 5;

 #pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
 #pragma omp target data map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
 #pragma omp target teams map(l)
   foo();
 
 #pragma omp target data map(always, tofrom: x)
 #pragma omp target data map(always: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 
 #pragma omp target teams private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target teams' directive}}  expected-note {{defined as private}}
   {}
 
 #pragma omp target teams firstprivate(j) map(j)  // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}} expected-note {{defined as firstprivate}}
   {}
 
 #pragma omp target teams map(m)
   {}
 
   return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
 }
 #endif