[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

4317

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

9346

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

9429

Same, too many params

lib/Sema/SemaLookup.cpp

1074 ↗

(On Diff #186338)

Why do we need special processing of the mapper here?

1793 ↗

(On Diff #186338)

Again, what's so special in mapper lookup?

lib/Sema/SemaOpenMP.cpp

13189–13195

I think this can be simplified:

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

13231–13237

Also can be simplified.

13358–13364

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

4317

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

lib/Sema/SemaLookup.cpp

1074 ↗

(On Diff #186338)

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 ↗

(On Diff #186338)

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

4317

Yes, it would be good

lib/Sema/SemaLookup.cpp

1074 ↗

(On Diff #186338)

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

4317

Then I'll do that. Thanks

lib/Sema/SemaLookup.cpp

1074 ↗

(On Diff #186338)

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 ↗

(On Diff #186338)

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 ↗

(On Diff #186338)

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 ↗

(On Diff #186338)

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 ↗

(On Diff #186338)

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 ↗

(On Diff #186338)

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

4201

assoicated -> associated

4310

assoicated -> associated

4325

amount -> number?

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

4234

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

4234

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

4234

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

4234

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

4234

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

3652

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

3682–3685

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.

4318

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

include/clang/Sema/Sema.h

9431

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

3682–3685

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

3670

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

9346

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/DeclOpenMP.h

 /// 'len' and 'data' should be mapped when mapping instances of 'struct vec'.
 ///
 /// \code
 /// #pragma omp declare mapper(mid: struct vec v) map(v.len, v.data[0:N])
 /// \endcode
 class OMPDeclareMapperDecl final : public ValueDecl, public DeclContext {
   friend class ASTDeclReader;
 
-  /// Clauses assoicated with this mapper declaration
+  /// Clauses associated with this mapper declaration
   MutableArrayRef<OMPClause *> Clauses;
 
   /// Mapper variable, which is 'v' in the example above
   Expr *MapperVarRef = nullptr;
 
   /// Name of the mapper variable
   DeclarationName VarName;
 

include/clang/AST/OpenMPClause.h

class OMPMappableExprListClause : public OMPVarListClause<T>,
   unsigned NumUniqueDeclarations;
 
   /// Number of component lists in this clause.
   unsigned NumComponentLists;
 
   /// Total number of components in this clause.
   unsigned NumComponents;
 
+  /// C++ nested name specifier for the associated user-defined mapper.
+  NestedNameSpecifierLoc MapperQualifierLoc;
+
+  /// The associated user-defined mapper identifier information.
+  DeclarationNameInfo MapperIdInfo;
+
+public:
+  /// This structure contains all sizes needed for by a
+  /// OMPMappableExprListClause.
+  struct OMPMappableExprListSizeTy {
+    unsigned NumVars;
+    unsigned NumUniqueDeclarations;
+    unsigned NumComponentLists;
+    unsigned NumComponents;
+  };
+
 protected:
   /// Build a clause for \a NumUniqueDeclarations declarations, \a
   /// 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)
+  /// \param MapperQualifierLocP C++ nested name specifier for the associated
+  /// user-defined mapper.
+  /// \param MapperIdInfoP The identifier of associated user-defined mapper.
+  OMPMappableExprListClause(
+      OpenMPClauseKind K, SourceLocation StartLoc, SourceLocation LParenLoc,
+      SourceLocation EndLoc, unsigned NumVars, unsigned NumUniqueDeclarations,
+      unsigned NumComponentLists, unsigned NumComponents,
+      NestedNameSpecifierLoc *MapperQualifierLocP = nullptr,
+      DeclarationNameInfo *MapperIdInfoP = nullptr)
       : OMPVarListClause<T>(K, StartLoc, LParenLoc, EndLoc, NumVars),
         NumUniqueDeclarations(NumUniqueDeclarations),
-        NumComponentLists(NumComponentLists), NumComponents(NumComponents) {}
+        NumComponentLists(NumComponentLists), NumComponents(NumComponents) {
+    if (MapperQualifierLocP)
+      MapperQualifierLoc = *MapperQualifierLocP;
+    if (MapperIdInfoP)
+      MapperIdInfo = *MapperIdInfoP;
+  }
+
+  /// Build a clause for \a NumUniqueDeclarations declarations, \a
+  /// 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 Sizes All required sizes to build a map clause. It includes 1)
+  /// NumVars: number of expressions listed in this clause; 2)
+  /// NumUniqueDeclarations: number of unique base declarations in this clause;
+  /// 3) NumComponentLists: number of component lists in this clause; and 4)
+  /// NumComponents: total number of expression components in the clause.
+  /// \param MapperQualifierLocP C++ nested name specifier for the associated
+  /// user-defined mapper.
+  /// \param MapperIdInfoP The identifier of associated user-defined mapper.
+  OMPMappableExprListClause(
+      OpenMPClauseKind K, SourceLocation StartLoc, SourceLocation LParenLoc,
+      SourceLocation EndLoc, OMPMappableExprListSizeTy Sizes,
+      NestedNameSpecifierLoc *MapperQualifierLocP = nullptr,
+      DeclarationNameInfo *MapperIdInfoP = nullptr)
+      : OMPVarListClause<T>(K, StartLoc, LParenLoc, EndLoc, Sizes.NumVars),
+        NumUniqueDeclarations(Sizes.NumUniqueDeclarations),
+        NumComponentLists(Sizes.NumComponentLists),
+        NumComponents(Sizes.NumComponents) {
+    if (MapperQualifierLocP)
+      MapperQualifierLoc = *MapperQualifierLocP;
+    if (MapperIdInfoP)
+      MapperIdInfo = *MapperIdInfoP;
+  }
 
   /// 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);
   }
for (auto &M : ComponentListMap) {
         ++CLSI;
 
         // Append components after the current components iterator.
         CI = std::copy(C.begin(), C.end(), CI);
       }
     }
   }
 
+  /// 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 *>(
+        static_cast<T *>(this)->template getTrailingObjects<Expr *>() +
+            OMPVarListClause<T>::varlist_size(),
+        OMPVarListClause<T>::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 *>(
+        static_cast<T *>(this)->template getTrailingObjects<Expr *>() +
+            OMPVarListClause<T>::varlist_size(),
+        OMPVarListClause<T>::varlist_size());
+  }
+
+  /// Set the user-defined mappers that are in the trailing objects of the
+  /// class.
+  void setUDMapperRefs(ArrayRef<Expr *> DMDs) {
+    assert(DMDs.size() == OMPVarListClause<T>::varlist_size() &&
+           "Unexpected number of user-defined mappers.");
+    std::copy(DMDs.begin(), DMDs.end(), getUDMapperRefs().begin());
+  }
+
 public:
   /// Return the number of unique base declarations in this clause.
   unsigned getUniqueDeclarationsNum() const { return NumUniqueDeclarations; }
 
   /// Return the number of lists derived from the clause expressions.
   unsigned getTotalComponentListNum() const { return NumComponentLists; }
 
   /// Return the total number of components in all lists derived from the
   /// clause.
   unsigned getTotalComponentsNum() const { return NumComponents; }
 
+  /// Gets the nested name specifier for associated user-defined mapper.
+  NestedNameSpecifierLoc getMapperQualifierLoc() const {
+    return MapperQualifierLoc;
+  }
+
+  /// Gets the name info for associated user-defined mapper.
+  const DeclarationNameInfo &getMapperIdInfo() const { return MapperIdInfo; }
+
   /// Iterator that browse the components by lists. It also allows
   /// browsing components of a single declaration.
   class const_component_lists_iterator
       : public llvm::iterator_adaptor_base<
             const_component_lists_iterator,
             MappableExprComponentListRef::const_iterator,
             std::forward_iterator_tag, MappableComponent, ptrdiff_t,
             MappableComponent, MappableComponent> {
public:
   using const_all_components_iterator = ArrayRef<MappableComponent>::iterator;
   using const_all_components_range =
       llvm::iterator_range<const_all_components_iterator>;
 
   const_all_components_range all_components() const {
     auto A = getComponentsRef();
     return const_all_components_range(A.begin(), A.end());
   }
+
+  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());
+  }
 };
 
 /// This represents clause 'map' in the '#pragma omp ...'
 /// directives.
 ///
 /// \code
 /// #pragma omp target map(a,b)
 /// \endcode
 /// In this example directive '#pragma omp target' has clause 'map'
 /// with the variables 'a' and 'b'.
 class OMPMapClause final : public OMPMappableExprListClause<OMPMapClause>,
                            private llvm::TrailingObjects<
                                OMPMapClause, Expr *, ValueDecl *, unsigned,
                                OMPClauseMappableExprCommon::MappableComponent> {
   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();
+    // There are varlist_size() of expressions, and varlist_size() of
+    // user-defined mappers.
+    return 2 * varlist_size();
   }
   size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
     return getUniqueDeclarationsNum();
   }
   size_t numTrailingObjects(OverloadToken<unsigned>) const {
     return getUniqueDeclarationsNum() + getTotalComponentListNum();
   }
 
 public:
   /// 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];
 
   /// 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 associated
+  /// 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.
+  /// \param Sizes All required sizes to build a map clause. It includes 1)
+  /// NumVars: number of expressions listed in this clause; 2)
+  /// NumUniqueDeclarations: number of unique base declarations in this clause;
+  /// 3) NumComponentLists: number of component lists in this clause; and 4)
+  /// NumComponents: total number of expression components in the clause.
   explicit 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));
+                        OMPMappableExprListSizeTy Sizes)
+      : OMPMappableExprListClause(OMPC_map, StartLoc, LParenLoc, EndLoc, Sizes,
+                                  &MapperQualifierLoc, &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.
-  /// \param NumComponents Total number of expression components in the clause.
-  explicit OMPMapClause(unsigned NumVars, unsigned NumUniqueDeclarations,
-                        unsigned NumComponentLists, unsigned NumComponents)
-      : OMPMappableExprListClause(
-            OMPC_map, SourceLocation(), SourceLocation(), SourceLocation(),
-            NumVars, NumUniqueDeclarations, NumComponentLists, NumComponents) {}
+  /// \param Sizes All required sizes to build a map clause. It includes 1)
+  /// NumVars: number of expressions listed in this clause; 2)
+  /// NumUniqueDeclarations: number of unique base declarations in this clause;
+  /// 3) NumComponentLists: number of component lists in this clause; and 4)
+  /// NumComponents: total number of expression components in the clause.
+  explicit OMPMapClause(OMPMappableExprListSizeTy Sizes)
+      : OMPMappableExprListClause(OMPC_map, SourceLocation(), SourceLocation(),
+                                  SourceLocation(), Sizes) {}
 
   /// Set map-type-modifier for the clause.
   ///
   /// \param I index for map-type-modifier.
   /// \param T map-type-modifier for the clause.
   void setMapTypeModifier(unsigned I, OpenMPMapModifierKind T) {
     assert(I < NumberOfModifiers &&
            "Unexpected index to store map type modifier, exceeds array size.");
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 associated
+  /// 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
-  /// clause.
-  /// \param NumComponentLists Number of unique base declarations in this
-  /// clause.
-  /// \param NumComponents Total number of expression components in the clause.
-  static OMPMapClause *CreateEmpty(const ASTContext &C, unsigned NumVars,
-                                   unsigned NumUniqueDeclarations,
-                                   unsigned NumComponentLists,
-                                   unsigned NumComponents);
+  /// \param Sizes All required sizes to build a map clause. It includes 1)
+  /// NumVars: number of expressions listed in this clause; 2)
+  /// NumUniqueDeclarations: number of unique base declarations in this clause;
+  /// 3) NumComponentLists: number of component lists in this clause; and 4)
+  /// NumComponents: total number of expression components in the clause.
+  static OMPMapClause *CreateEmpty(const ASTContext &C,
+                                   OMPMappableExprListSizeTy Sizes);
 
   /// Fetches mapping kind for the clause.
   OpenMPMapClauseKind getMapType() const LLVM_READONLY { return MapType; }
 
   /// Is this an implicit map type?
   /// We have to capture 'IsMapTypeImplicit' from the parser for more
   /// informative error messages.  It helps distinguish map(r) from
   /// map(tofrom: r), which is important to print more helpful error
   /// messages for some target directives.
   bool isImplicitMapType() const LLVM_READONLY { return MapTypeIsImplicit; }
 
   /// Fetches the map-type-modifier at 'Cnt' index of array of modifiers.
   ///
   /// \param Cnt index for map-type-modifier.
   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);
   }
 
   /// Fetches location of clause mapping kind.
   SourceLocation getMapLoc() const LLVM_READONLY { return MapLoc; }
 
class OMPIsDevicePtrClause final
 
   /// 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.
   /// \param NumComponents Total number of expression components in the clause.
+  /// \param Sizes All required sizes to build a map clause. It includes 1)
+  /// NumVars: number of expressions listed in this clause; 2)
+  /// NumUniqueDeclarations: number of unique base declarations in this clause;
+  /// 3) NumComponentLists: number of component lists in this clause; and 4)
+  /// NumComponents: total number of expression components in the clause.
   explicit OMPIsDevicePtrClause(unsigned NumVars,
                                 unsigned NumUniqueDeclarations,
                                 unsigned NumComponentLists,
                                 unsigned NumComponents)
       : OMPMappableExprListClause(OMPC_is_device_ptr, SourceLocation(),
                                   SourceLocation(), SourceLocation(), NumVars,
                                   NumUniqueDeclarations, NumComponentLists,
                                   NumComponents) {}

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

public:
   /// \param RLoc Returned location of right paren.
   ExprResult ParseOpenMPParensExpr(StringRef ClauseName, SourceLocation &RLoc);
 
   /// Data used for parsing list of variables in OpenMP clauses.
   struct OpenMPVarListDataTy {
     Expr *TailExpr = nullptr;
     SourceLocation ColonLoc;
     SourceLocation RLoc;
-    CXXScopeSpec ReductionIdScopeSpec;
-    DeclarationNameInfo ReductionId;
+    CXXScopeSpec ReductionOrMapperIdScopeSpec;
+    DeclarationNameInfo ReductionOrMapperId;
     OpenMPDependClauseKind DepKind = OMPC_DEPEND_unknown;
     OpenMPLinearClauseKind LinKind = OMPC_LINEAR_val;
     SmallVector<OpenMPMapModifierKind, OMPMapClause::NumberOfModifiers>
     MapTypeModifiers;
     SmallVector<SourceLocation, OMPMapClause::NumberOfModifiers>
     MapTypeModifiersLoc;
     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

public:
   OMPClause *ActOnOpenMPAtomicDefaultMemOrderClause(
       OpenMPAtomicDefaultMemOrderClauseKind Kind, SourceLocation KindLoc,
       SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc);
 
   OMPClause *ActOnOpenMPVarListClause(
       OpenMPClauseKind Kind, ArrayRef<Expr *> Vars, Expr *TailExpr,
       SourceLocation StartLoc, SourceLocation LParenLoc,
       SourceLocation ColonLoc, SourceLocation EndLoc,
-      CXXScopeSpec &ReductionIdScopeSpec,
-      const DeclarationNameInfo &ReductionId, OpenMPDependClauseKind DepKind,
+      CXXScopeSpec &ReductionOrMapperIdScopeSpec,
+      DeclarationNameInfo &ReductionOrMapperId, OpenMPDependClauseKind DepKind,
       OpenMPLinearClauseKind LinKind,
       ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
-      ArrayRef<SourceLocation> MapTypeModifiersLoc,
-      OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
-      SourceLocation DepLinMapLoc);
+      ArrayRef<SourceLocation> MapTypeModifiersLoc, 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.
   OMPClause *ActOnOpenMPFirstprivateClause(ArrayRef<Expr *> VarList,
                                            SourceLocation StartLoc,
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) {
-  unsigned NumVars = Vars.size();
-  unsigned NumUniqueDeclarations =
-      getUniqueDeclarationsTotalNumber(Declarations);
-  unsigned NumComponentLists = ComponentLists.size();
-  unsigned NumComponents = getComponentsTotalNumber(ComponentLists);
+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) {
+  OMPMappableExprListSizeTy Sizes;
+  Sizes.NumVars = Vars.size();
+  Sizes.NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations);
+  Sizes.NumComponentLists = ComponentLists.size();
+  Sizes.NumComponents = getComponentsTotalNumber(ComponentLists);
 
   // We need to allocate:
-  // NumVars x Expr* - we have an original list expression for each clause list
-  // entry.
+  // 2 x NumVars x Expr* - we have an original list expression and an associated
+  // user-defined mapper for each clause list entry.
   // NumUniqueDeclarations x ValueDecl* - unique base declarations associated
   // with each component list.
   // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
   // number of lists for each unique declaration and the size of each component
   // list.
   // NumComponents x MappableComponent - the total of all the components in all
   // the lists.
   void *Mem = C.Allocate(
       totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
-          NumVars, NumUniqueDeclarations,
-          NumUniqueDeclarations + NumComponentLists, NumComponents));
+          2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
+          Sizes.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, Sizes);
 
   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) {
+OMPMapClause *OMPMapClause::CreateEmpty(const ASTContext &C,
+                                        OMPMappableExprListSizeTy Sizes) {
   void *Mem = C.Allocate(
       totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
-          NumVars, NumUniqueDeclarations,
-          NumUniqueDeclarations + NumComponentLists, NumComponents));
-  return new (Mem) OMPMapClause(NumVars, NumUniqueDeclarations,
-                                NumComponentLists, NumComponents);
+          2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
+          Sizes.NumComponents));
+  return new (Mem) OMPMapClause(Sizes);
 }
 
 OMPToClause *OMPToClause::Create(const ASTContext &C, SourceLocation StartLoc,
                                  SourceLocation LParenLoc,
                                  SourceLocation EndLoc, ArrayRef<Expr *> Vars,
                                  ArrayRef<ValueDecl *> Declarations,
                                  MappableExprComponentListsRef ComponentLists) {
   unsigned NumVars = Vars.size();

 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.ReductionOrMapperIdScopeSpec,
+                                       /*ObjectType=*/nullptr,
+                                       /*EnteringContext=*/false);
+      if (Tok.isNot(tok::identifier) && Tok.isNot(tok::kw_default)) {
+        Diag(Tok.getLocation(), diag::err_omp_mapper_illegal_identifier);
+        SkipUntil(tok::colon, tok::annot_pragma_openmp_end, StopBeforeMatch);
+        return true;
+      }
+      auto &DeclNames = Actions.getASTContext().DeclarationNames;
+      Data.ReductionOrMapperId = DeclarationNameInfo(
+          DeclNames.getIdentifier(Tok.getIdentifierInfo()), Tok.getLocation());
+      ConsumeToken();
+      // Parse ')'.
+      T.consumeClose();
     } else {
       // For the case of unknown map-type-modifier or a map-type.
       // Map-type is followed by a colon; the function returns when it
       // 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;
   BalancedDelimiterTracker LinearT(*this, tok::l_paren,
                                   tok::annot_pragma_openmp_end);
   // Handle reduction-identifier for reduction clause.
   if (Kind == OMPC_reduction || Kind == OMPC_task_reduction ||
       Kind == OMPC_in_reduction) {
     ColonProtectionRAIIObject ColonRAII(*this);
     if (getLangOpts().CPlusPlus)
-      ParseOptionalCXXScopeSpecifier(Data.ReductionIdScopeSpec,
+      ParseOptionalCXXScopeSpecifier(Data.ReductionOrMapperIdScopeSpec,
                                      /*ObjectType=*/nullptr,
                                      /*EnteringContext=*/false);
-    InvalidReductionId = ParseReductionId(*this, Data.ReductionIdScopeSpec,
-                                          UnqualifiedReductionId);
+    InvalidReductionId = ParseReductionId(
+        *this, Data.ReductionOrMapperIdScopeSpec, UnqualifiedReductionId);
     if (InvalidReductionId) {
       SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end,
                 StopBeforeMatch);
     }
     if (Tok.is(tok::colon))
       Data.ColonLoc = ConsumeToken();
     else
       Diag(Tok, diag::warn_pragma_expected_colon) << "reduction identifier";
     if (!InvalidReductionId)
-      Data.ReductionId =
+      Data.ReductionOrMapperId =
           Actions.GetNameFromUnqualifiedId(UnqualifiedReductionId);
   } else if (Kind == OMPC_depend) {
   // Handle dependency type for depend clause.
     ColonProtectionRAIIObject ColonRAII(*this);
     Data.DepKind =
         static_cast<OpenMPDependClauseKind>(getOpenMPSimpleClauseType(
             Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : ""));
     Data.DepLinMapLoc = Tok.getLocation();
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.ReductionOrMapperIdScopeSpec, Data.ReductionOrMapperId, Data.DepKind,
+      Data.LinKind, Data.MapTypeModifiers, Data.MapTypeModifiersLoc,
+      Data.MapType, Data.IsMapTypeImplicit, Data.DepLinMapLoc);
 }
 

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::ActOnOpenMPDynamicAllocatorsClause(SourceLocation StartLoc,
                                                     SourceLocation EndLoc) {
   return new (Context) OMPDynamicAllocatorsClause(StartLoc, EndLoc);
 }
 
 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,
+    SourceLocation EndLoc, CXXScopeSpec &ReductionOrMapperIdScopeSpec,
+    DeclarationNameInfo &ReductionOrMapperId, OpenMPDependClauseKind DepKind,
     OpenMPLinearClauseKind LinKind,
     ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
-    ArrayRef<SourceLocation> MapTypeModifiersLoc,
-    OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
-    SourceLocation DepLinMapLoc) {
+    ArrayRef<SourceLocation> MapTypeModifiersLoc, OpenMPMapClauseKind MapType,
+    bool IsMapTypeImplicit, SourceLocation DepLinMapLoc) {
   OMPClause *Res = nullptr;
   switch (Kind) {
   case OMPC_private:
     Res = ActOnOpenMPPrivateClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_firstprivate:
     Res = ActOnOpenMPFirstprivateClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_lastprivate:
     Res = ActOnOpenMPLastprivateClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_shared:
     Res = ActOnOpenMPSharedClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_reduction:
     Res = ActOnOpenMPReductionClause(VarList, StartLoc, LParenLoc, ColonLoc,
-                                     EndLoc, ReductionIdScopeSpec, ReductionId);
+                                     EndLoc, ReductionOrMapperIdScopeSpec,
+                                     ReductionOrMapperId);
     break;
   case OMPC_task_reduction:
     Res = ActOnOpenMPTaskReductionClause(VarList, StartLoc, LParenLoc, ColonLoc,
-                                         EndLoc, ReductionIdScopeSpec,
-                                         ReductionId);
+                                         EndLoc, ReductionOrMapperIdScopeSpec,
+                                         ReductionOrMapperId);
     break;
   case OMPC_in_reduction:
-    Res =
-        ActOnOpenMPInReductionClause(VarList, StartLoc, LParenLoc, ColonLoc,
-                                     EndLoc, ReductionIdScopeSpec, ReductionId);
+    Res = ActOnOpenMPInReductionClause(VarList, StartLoc, LParenLoc, ColonLoc,
+                                       EndLoc, ReductionOrMapperIdScopeSpec,
+                                       ReductionOrMapperId);
     break;
   case OMPC_linear:
     Res = ActOnOpenMPLinearClause(VarList, TailExpr, StartLoc, LParenLoc,
                                   LinKind, DepLinMapLoc, ColonLoc, EndLoc);
     break;
   case OMPC_aligned:
     Res = ActOnOpenMPAlignedClause(VarList, TailExpr, StartLoc, LParenLoc,
                                    ColonLoc, EndLoc);
     break;
   case OMPC_copyin:
     Res = ActOnOpenMPCopyinClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_copyprivate:
     Res = ActOnOpenMPCopyprivateClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   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,
-                               IsMapTypeImplicit, DepLinMapLoc, ColonLoc,
-                               VarList, StartLoc, LParenLoc, EndLoc);
+    Res = ActOnOpenMPMapClause(
+        MapTypeModifiers, MapTypeModifiersLoc, ReductionOrMapperIdScopeSpec,
+        ReductionOrMapperId, MapType, IsMapTypeImplicit, DepLinMapLoc, ColonLoc,
     break;
   case OMPC_to:
     Res = ActOnOpenMPToClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_from:
     Res = ActOnOpenMPFromClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_use_device_ptr:
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 (auto RD : D->redecls()) {
     if (LookupResult::isVisible(SemaRef, ND))
       return ND;
   }
 
   return nullptr;
 }
 
 static void
-argumentDependentLookup(Sema &SemaRef, const DeclarationNameInfo &ReductionId,
+argumentDependentLookup(Sema &SemaRef, const DeclarationNameInfo &Id,
                         SourceLocation Loc, QualType Ty,
                         SmallVectorImpl<UnresolvedSet<8>> &Lookups) {
   // Find all of the associated namespaces and classes based on the
   // arguments we have.
   Sema::AssociatedNamespaceSet AssociatedNamespaces;
   Sema::AssociatedClassSet AssociatedClasses;
   OpaqueValueExpr OVE(Loc, Ty, VK_LValue);
   SemaRef.FindAssociatedClassesAndNamespaces(Loc, &OVE, AssociatedNamespaces,
for (auto *NS : AssociatedNamespaces) {
     //
     //     -- Any using-directives in the associated namespace are
     //        ignored.
     //
     //     -- Any namespace-scope friend functions declared in
     //        associated classes are visible within their respective
     //        namespaces even if they are not visible during an ordinary
     //        lookup (11.4).
-    DeclContext::lookup_result R = NS->lookup(ReductionId.getName());
+    DeclContext::lookup_result R = NS->lookup(Id.getName());
     for (auto *D : R) {
       auto *Underlying = D;
       if (auto *USD = dyn_cast<UsingShadowDecl>(D))
         Underlying = USD->getTargetDecl();
 
-      if (!isa<OMPDeclareReductionDecl>(Underlying))
+      if (!isa<OMPDeclareReductionDecl>(Underlying) &&
+          !isa<OMPDeclareMapperDecl>(Underlying))
         continue;
 
       if (!SemaRef.isVisible(D)) {
         D = findAcceptableDecl(SemaRef, D);
         if (!D)
           continue;
         if (auto *USD = dyn_cast<UsingShadowDecl>(D))
           Underlying = USD->getTargetDecl();
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,
+                                     const 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();
+  SourceLocation Loc = MapperId.getLoc();
+  // Perform argument dependent lookup.
+  if (SemaRef.getLangOpts().CPlusPlus && !MapperIdScopeSpec.isSet())
+    argumentDependentLookup(SemaRef, MapperId, Loc, Type, Lookups);
+  // Return the first user-defined mapper with the desired type.
+  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,
+    const 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++;
+      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;
+        MVLI.UDMapperList.push_back(ER.get());
+      } 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 (!RE->IgnoreParenImpCasts()->isLValue()) {
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;
+        MVLI.UDMapperList.push_back(ER.get());
+      } 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;
+      MVLI.UDMapperList.push_back(ER.get());
+    } 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() || 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

case OMPC_depend: {
     unsigned NumLoops = Record.readInt();
     C = OMPDependClause::CreateEmpty(Context, NumVars, NumLoops);
     break;
   }
   case OMPC_device:
     C = new (Context) OMPDeviceClause();
     break;
   case OMPC_map: {
-    unsigned NumVars = Record.readInt();
-    unsigned NumDeclarations = Record.readInt();
-    unsigned NumLists = Record.readInt();
-    unsigned NumComponents = Record.readInt();
-    C = OMPMapClause::CreateEmpty(Context, NumVars, NumDeclarations, NumLists,
-                                  NumComponents);
+    OMPMapClause::OMPMappableExprListSizeTy Sizes;
+    Sizes.NumVars = Record.readInt();
+    Sizes.NumUniqueDeclarations = Record.readInt();
+    Sizes.NumComponentLists = Record.readInt();
+    Sizes.NumComponents = Record.readInt();
+    C = OMPMapClause::CreateEmpty(Context, Sizes);
     break;
   }
   case OMPC_num_teams:
     C = new (Context) OMPNumTeamsClause();
     break;
   case OMPC_thread_limit:
     C = new (Context) OMPThreadLimitClause();
     break;

 
 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_codegen.cpp

+///==========================================================================///
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY0 %s
+
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+class C {
+public:
+  int a;
+};
+
+#pragma omp declare mapper(id: C s) map(s.a)
+
+// CHECK-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l49.region_id = weak constant i8 0
+
+// CHECK-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i{{64|32}}] [i{{64|32}} 4]
+// CHECK-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35]
+
+// CHECK-LABEL: foo{{.*}}(
+void foo(int a){
+  int i = a;
+  C c;
+  c.a = a;
+
+  // CHECK-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
+  // CHECK-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
+  // CHECK-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
+  // CHECK-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
+  // CHECK-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
+  // CHECK-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
+  // CHECK-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
+  // CHECK-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
+  // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
+  #pragma omp target map(mapper(default),tofrom: c)
+  {
+   ++c.a;
+  }
+  // CHECK: call void [[KERNEL:@.+]](%class.C* [[VAL]])
+}
+
+
+// CHECK: define internal void [[KERNEL]](%class.C* {{.+}}[[ARG:%.+]])
+// CHECK: [[ADDR:%.+]] = alloca %class.C*,
+// CHECK: store %class.C* [[ARG]], %class.C** [[ADDR]]
+// CHECK: [[CADDR:%.+]] = load %class.C*, %class.C** [[ADDR]]
+// CHECK: [[CAADDR:%.+]] = getelementptr inbounds %class.C, %class.C* [[CADDR]], i32 0, i32 0
+// CHECK: [[VAL:%[^,]+]] = load i32, i32* [[CAADDR]]
+// CHECK: {{.+}} = add nsw i32 [[VAL]], 1
+// CHECK: }
+
+#endif

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