[OpenMP 5.0] Codegen support to pass user-defined mapper functions to runtime

[lildmh]

Summary

This patch implements the code generation to use OpenMP 5.0 declare mapper (e.g., user-defined mapper) constructs. It looks up the proper mapper function for each map, to, or from clause that has a user-defined mapper associated, and passes them to the OpenMP runtime function.
The design slides can be found at https://github.com/lingda-li/public-sharing/blob/master/mapper_runtime_design.pptx

[lildmh]

Rebase

[lildmh]

Rebase

[ABataev]

I would suggest to split this patch into 2 part: the first one is NFC and just makes the ompiler to use the new interfaces and the second with the new functionality for mappers.

include/clang/AST/OpenMPClause.h

4292

Not in Camel format

4537–4541

Why changed to const Expr * here? I believe, ArrayRef makes it constant automatically.

4590

Not in Camel format

4593

No need for const in ArrayRef

4615–4616

Not in Camel format

4737

ArrayRef<const Expr *>()->llvm::None?

4744

ArrayRef<const Expr *>()->llvm::None

4757

Same here

lib/CodeGen/CGOpenMPRuntime.cpp

743

Better to fix it in a separate patch

756–795

I believe these new functions replace completely the old one, so old functions must be removed

2688

Same here, remove processing for old functions.

7440

CamelCase

7440

Seems to me, with mapper you're changing the contract of this function. It is supposed to return the size in bytes, with Mapper it returns just number of elements. Bad decision. Better to convert size in bytes to number of elements in place where you need it.

7461

Why do we return 1 here?

7462

No need for else here

7481–7484

Same here

7488–7497

Same here

7489

No need for else here

7592–7601

Too many params in the function, worth thinking about wrapping them in a record.

7968–7971

Mappers.push_back(HasMappers?Mapper:nullptr);

8751–8753

const auto *VD = dyn_cast_or_null<VarDecl>(std::get<0>(L));

8856–8857

In general, there should not be branch for default mapping during the codegen, the implicit map clauses must be generated in sema.

8988–8990

Better to use Builder.CreateConstArrayGEP

8991–8992

Easier cast function to the voidptr, I think

9024–9027

Hmm, just cast of the Info.MappersArray to the correct type if all indices are 0?

9028–9029

If then sub-statement in braces, the else substatement also must be in braces

9385

Better to use MapperCGF.EmitNounwindRuntimeCall.

lib/CodeGen/CGOpenMPRuntime.h

1519

Camel case

1533

The flag HasMappers also must be cleared, no?

[lildmh]

Rebase and address comments

[lildmh]

I'll split the patch into 2 later

lib/CodeGen/CGOpenMPRuntime.cpp

7440

Yes. The reason why I did this is I found this is the single place that I can consider all situations. Otherwise, I need to divide the size by the element size in other places. Since this function has discussed all situations for the element size, I will need to duplicate the work elsewhere if I don't do it here. The latter solution is not good for code mountainous I think.

7592–7601

Maybe not in this patch?

8856–8857

Not sure where this code will be used. I guess it's still correct to add a null to the mappers here?

8988–8990

I saw all the other places are using CreateConstInBoundsGEP2_32 to get sizes, types, etc. So I kept it consistent here. Maybe these modifications should be in a future patch?

8991–8992

Again, this is to keep it consistent with previous code above?

9024–9027

Again, this is to keep consistent with previous code?

[ABataev]

clang/include/clang/AST/OpenMPClause.h

4918 ↗

(On Diff #234113)

Do we really need to set HasMapper to true unconditionally here and in other places?

clang/lib/CodeGen/CGOpenMPRuntime.h

841–842 ↗

(On Diff #234113)

Is this required to make it virtual?

1556 ↗

(On Diff #234113)

(!HasMapper || MappersArray)?

lib/CodeGen/CGOpenMPRuntime.cpp

7440

Changing the original function contract is also not the best solution. Better to introduce the new function. If some of the code can be reused, need to factor out the common code and reuse it in the old function and the new one.

7440

Also, additional question. Does it mean that we pass to the runtime functions different values in different situations: in one case we pass size in bytes, with mappers - number of elements? If so, it is a very bad idea caused by bad design. We need to pass either size in bytes, or number of elements in this argument.

7592–7601

It is better to make it in this patch or make a separate patch and commit it before this one. We try to resolve this as soon as possible.

8856–8857

I mean, that in general, it would be good to see here some kind of an assertion. Almost all situations must be resolved in Sema.

8988–8990

Some of the modifications were made before we had these new functions. Now, we have the new ones and if we can use it, better to switch to the new safer versions.

8991–8992

No need to rely on the old code if can make something cleaner and better. Cast Mfunc to voidptr in the if statement and no need to make some extra casts later.

9024–9027

Same, use the best code if we can use it.

[lildmh]

Alexey, thanks for the review

clang/include/clang/AST/OpenMPClause.h

4918 ↗

(On Diff #234113)

For map, to, and from clauses, they can have mappers, so it's set to true. For is_device_ptr and use_device_ptr, it's set to false.

lib/CodeGen/CGOpenMPRuntime.cpp

7440

Yes, the current design is like this. Let me see how to make it better

8856–8857

Okay, I guess it's not the work of this patch

[ABataev]

clang/include/clang/AST/OpenMPClause.h

4918 ↗

(On Diff #234113)

So, it completely depends on the clause kind, right? If so, can we just remove it and rely on the clause kind?

[lildmh]

clang/include/clang/AST/OpenMPClause.h

4918 ↗

(On Diff #234113)

This is used in OMPMappableExprListClause which is inherited by all these clauses. Do you mean to check the template type there to get the clause kind?

[ABataev]

clang/include/clang/AST/OpenMPClause.h

4918 ↗

(On Diff #234113)

I see. I think it is better to rename it to ClauseSupportsMapper (or something like this) and make it const. Currently, it is confusing. I thought that this flag shows the use of the mapper in the clause.

[lildmh]

clang/include/clang/AST/OpenMPClause.h

4918 ↗

(On Diff #234113)

Ok, will do that

[lildmh]

Address Alexey's comments to change mapper function size and refactor code

[ABataev]

Looks much better in general, need to resolve last one issue with number of elements/size and you're ready.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8981–8982 ↗

(On Diff #234554)

So, we're still going to use number of elements for mappers? And pass it in the same parameter that in other cases is used as size in bytes? If so, point to it explicitly in the review for the runtime part so all are informed about it.

9244 ↗

(On Diff #234554)

getOrEmitUserDefinedMapperFunc?

[lildmh]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8981–8982 ↗

(On Diff #234554)

From interface, the mapper function uses size in bytes now. Inside, it needs number of elements to iterate through all elements. This has no impact on the runtime part, since it looks like normal mapping from the interface. All conversion happens inside the mapper function which is completely generated by the compiler.

9244 ↗

(On Diff #234554)

I guess getUserDefinedMapperFunc is a better name? Because the user uses this function to get the mapper function. Emitting a mapper function is a side effect.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8981–8982 ↗

(On Diff #234554)

Ok. Then why do we need to convert size in bytes to number of elements here?

9244 ↗

(On Diff #234554)

I meant something like getOrCreate... but it is up to you.

[lildmh]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8981–8982 ↗

(On Diff #234554)

This is used to 1) see if we are going to map an array of elements with mapper, and 2) iterate all to map them individually.

9244 ↗

(On Diff #234554)

Okay, sounds good

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8981–8982 ↗

(On Diff #234554)

Could you point where we have this kind of analysis here? Because I don't see anything affected by this change in the patch.

[lildmh]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8986 ↗

(On Diff #234554)

Size is used to compute the last element position, which is used as the loop boundary.

9207 ↗

(On Diff #234554)

Size is used to see if this is an array.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8981–8982 ↗

(On Diff #234554)

Is this a bug fix in the previous implementation?

[lildmh]

Change the function name

[lildmh]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8981–8982 ↗

(On Diff #234554)

The previous implementation assumes the size is the number of elements, and it works correctly under that assumption. Since we change the meaning of size here, I add this line of code so the previous implementation can work correctly in the new assumption that the size is the size in bytes.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8981–8982 ↗

(On Diff #234554)

Ah, got it. Then, in general, looks good. Please, split the patches in to, possibly, 2 NFC (one with using of the new functions + another one for aggregating too many params into records) + another one with the new functionality.

[lildmh]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8981–8982 ↗

(On Diff #234554)

Okay, will do that. What do you think need to be done for the runtime patch D68100?

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8981–8982 ↗

(On Diff #234554)

Address comments and rebase, I think

[lildmh]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8981–8982 ↗

(On Diff #234554)

The only left comment is to split it into 2 patches, I think.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8981–8982 ↗

(On Diff #234554)

Do it.

[lildmh]

I'll update the diff and please check and accept after that

[ABataev]

In D67833#2151501, @lildmh wrote:

I'll update the diff and please check and accept after that

Ok

[lildmh]

Update diff and pass test. Please accept

[ABataev]

LG

[lildmh]

@grokos Could you upstream this patch and the runtime patch neck to neck? Upstreaming one of them will break the OpenMP offloading. It will be nice if you can test this patch locally with the runtime patch. Thanks

[grokos]

I tried to build clang with this patch and I get errors like:

CGOpenMPRuntime.cpp:9463:38: error: ‘OMPRTL___tgt_target_teams_nowait_mapper’ was not declared in this scope
                                    ? OMPRTL___tgt_target_teams_nowait_mapper

Where are these OMPRTL___tgt_ symbols defined?

[lildmh]

Include the llvm part

[lildmh]

In D67833#2154312, @grokos wrote:

I tried to build clang with this patch and I get errors like:

CGOpenMPRuntime.cpp:9463:38: error: ‘OMPRTL___tgt_target_teams_nowait_mapper’ was not declared in this scope
                                    ? OMPRTL___tgt_target_teams_nowait_mapper

Where are these OMPRTL___tgt_ symbols defined?

Sorry I forgot to include the llvm part. Please try again

[grokos]

OK, now it works. Thanks!

[fhahn]

I think this change introduced the following warnings, where auto & is used for types that are always copied. It would be great if you could take a look.

lvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8011:24: warning: loop variable 'L' is always a copy because the range of type 'clang::OMPMappableExprListClause<clang::OMPToClause>::const_component_lists_range' (aka 'iterator_range<clang::OMPMappableExprListClause<clang::OMPToClause>::const_component_lists_iterator>') does not return a reference [-Wrange-loop-analysis]
      for (const auto &L : C->component_lists()) {
                       ^
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8011:12: note: use non-reference type 'std::__1::tuple<const clang::ValueDecl *, llvm::ArrayRef<clang::OMPClauseMappableExprCommon::MappableComponent>, const clang::ValueDecl *>'
      for (const auto &L : C->component_lists()) {
           ^~~~~~~~~~~~~~~
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8016:24: warning: loop variable 'L' is always a copy because the range of type 'clang::OMPMappableExprListClause<clang::OMPFromClause>::const_component_lists_range' (aka 'iterator_range<clang::OMPMappableExprListClause<clang::OMPFromClause>::const_component_lists_iterator>') does not return a reference [-Wrange-loop-analysis]
      for (const auto &L : C->component_lists()) {
                       ^
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8016:12: note: use non-reference type 'std::__1::tuple<const clang::ValueDecl *, llvm::ArrayRef<clang::OMPClauseMappableExprCommon::MappableComponent>, const clang::ValueDecl *>'
      for (const auto &L : C->component_lists()) {
           ^~~~~~~~~~~~~~~
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8032:24: warning: loop variable 'L' is always a copy because the range of type 'clang::OMPMappableExprListClause<clang::OMPUseDevicePtrClause>::const_component_lists_range' (aka 'iterator_range<clang::OMPMappableExprListClause<clang::OMPUseDevicePtrClause>::const_component_lists_iterator>') does not return a reference [-Wrange-loop-analysis]
      for (const auto &L : C->component_lists()) {
                       ^
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8032:12: note: use non-reference type 'std::__1::tuple<const clang::ValueDecl *, llvm::ArrayRef<clang::OMPClauseMappableExprCommon::MappableComponent>, const clang::ValueDecl *>'
      for (const auto &L : C->component_lists()) {
           ^~~~~~~~~~~~~~~
3 warnings generated.

[lildmh]

In D67833#2155196, @fhahn wrote:

I think this change introduced the following warnings, where auto & is used for types that are always copied. It would be great if you could take a look.

lvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8011:24: warning: loop variable 'L' is always a copy because the range of type 'clang::OMPMappableExprListClause<clang::OMPToClause>::const_component_lists_range' (aka 'iterator_range<clang::OMPMappableExprListClause<clang::OMPToClause>::const_component_lists_iterator>') does not return a reference [-Wrange-loop-analysis]
      for (const auto &L : C->component_lists()) {
                       ^
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8011:12: note: use non-reference type 'std::__1::tuple<const clang::ValueDecl *, llvm::ArrayRef<clang::OMPClauseMappableExprCommon::MappableComponent>, const clang::ValueDecl *>'
      for (const auto &L : C->component_lists()) {
           ^~~~~~~~~~~~~~~
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8016:24: warning: loop variable 'L' is always a copy because the range of type 'clang::OMPMappableExprListClause<clang::OMPFromClause>::const_component_lists_range' (aka 'iterator_range<clang::OMPMappableExprListClause<clang::OMPFromClause>::const_component_lists_iterator>') does not return a reference [-Wrange-loop-analysis]
      for (const auto &L : C->component_lists()) {
                       ^
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8016:12: note: use non-reference type 'std::__1::tuple<const clang::ValueDecl *, llvm::ArrayRef<clang::OMPClauseMappableExprCommon::MappableComponent>, const clang::ValueDecl *>'
      for (const auto &L : C->component_lists()) {
           ^~~~~~~~~~~~~~~
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8032:24: warning: loop variable 'L' is always a copy because the range of type 'clang::OMPMappableExprListClause<clang::OMPUseDevicePtrClause>::const_component_lists_range' (aka 'iterator_range<clang::OMPMappableExprListClause<clang::OMPUseDevicePtrClause>::const_component_lists_iterator>') does not return a reference [-Wrange-loop-analysis]
      for (const auto &L : C->component_lists()) {
                       ^
llvm-project/clang/lib/CodeGen/CGOpenMPRuntime.cpp:8032:12: note: use non-reference type 'std::__1::tuple<const clang::ValueDecl *, llvm::ArrayRef<clang::OMPClauseMappableExprCommon::MappableComponent>, const clang::ValueDecl *>'
      for (const auto &L : C->component_lists()) {
           ^~~~~~~~~~~~~~~
3 warnings generated.

Thanks, I'll fix this soon. Basically &L needs to be L instead in those 3 places. Since I don't have upload permission, it will be great if you can do the change quickly

[lildmh]

Fix is at https://reviews.llvm.org/D83959
@Meinersbur Michael, could you please help upstream the patch above? Thanks

[Meinersbur]

Seems that it already has been applied ;-)

[lildmh]

In D67833#2161241, @Meinersbur wrote:

Seems that it already has been applied ;-)

Yes, thanks Michael. Hope everything goes well with you

[protze.joachim]

Starting with the commit of this patch, the libomptarget test llvm-project/openmp/libomptarget/test/offloading/target_depend_nowait.cpp fails. Here is a stacktrace of the segfault:

$ gdb target_depend_nowait.cpp.tmp-x86_64-pc-linux-gnu
(gdb) run
...
Program received signal SIGSEGV, Segmentation fault.
(gdb) bt
#0  0x0000000000400fcf in .omp_outlined._debug__ (.global_tid.=0x2aaab96b9be0, .bound_tid.=0x2aaab96b9bd8) at target_depend_nowait.cpp:22
#1  0x0000000000401e8d in .omp_outlined..23 (.global_tid.=0x2aaab96b9be0, .bound_tid.=0x2aaab96b9bd8) at target_depend_nowait.cpp:18
#2  0x00002aaaab574213 in __kmp_invoke_microtask () from libomp.so
#3  0x00002aaaab51338e in __kmp_invoke_task_func () from libomp.so
#4  0x00002aaaab5126bf in __kmp_launch_thread () from libomp.so
#5  0x00002aaaab55d3d0 in __kmp_launch_worker(void*) () from libomp.so
#6  0x00002aaaabbd6e65 in start_thread () from libpthread.so.0
#7  0x00002aaaabee988d in clone () from libc.so.6

[lildmh]

In D67833#2166478, @protze.joachim wrote:

Starting with the commit of this patch, the libomptarget test llvm-project/openmp/libomptarget/test/offloading/target_depend_nowait.cpp fails. Here is a stacktrace of the segfault:

$ gdb target_depend_nowait.cpp.tmp-x86_64-pc-linux-gnu
(gdb) run
...
Program received signal SIGSEGV, Segmentation fault.
(gdb) bt
#0  0x0000000000400fcf in .omp_outlined._debug__ (.global_tid.=0x2aaab96b9be0, .bound_tid.=0x2aaab96b9bd8) at target_depend_nowait.cpp:22
#1  0x0000000000401e8d in .omp_outlined..23 (.global_tid.=0x2aaab96b9be0, .bound_tid.=0x2aaab96b9bd8) at target_depend_nowait.cpp:18
#2  0x00002aaaab574213 in __kmp_invoke_microtask () from libomp.so
#3  0x00002aaaab51338e in __kmp_invoke_task_func () from libomp.so
#4  0x00002aaaab5126bf in __kmp_launch_thread () from libomp.so
#5  0x00002aaaab55d3d0 in __kmp_launch_worker(void*) () from libomp.so
#6  0x00002aaaabbd6e65 in start_thread () from libpthread.so.0
#7  0x00002aaaabee988d in clone () from libc.so.6

Thanks. If possible, could you send me an IR generated before this patch and one after this patch, so I can probably see the problem?

[lildmh]

I couldn't reproduce this problem because all commands are ignored on my machine. Could you send the exact compiling and running commands?
In the IR files you sent me before, any idea where segfault happens?

[protze.joachim]

I'm executing:

$ /usr/bin/python bin/llvm-lit -vv -a projects/openmp/libomptarget/test/offloading/target_depend_nowait.cpp

which executes:

"./bin/clang++" "-fopenmp" "-pthread" "-fno-experimental-isel" "-I" "../openmp/libomptarget/test" "-I" "./projects/openmp/libomptarget/../runtime/src" "-L" "./lib" "-fopenmp-targets=x86_64-pc-linux-gnu" "../openmp/libomptarget/test/offloading/target_depend_nowait.cpp" "-o" "./projects/openmp/libomptarget/test/offloading/Output/target_depend_nowait.cpp.tmp-x86_64-pc-linux-gnu"

According to gdb, the segfault happens in 0x400fcf. %rcx is 0 at that time.

 400f89:       e8 62 fc ff ff          callq  400bf0 <__kmpc_omp_target_task_alloc@plt>
  400f8e:       45 31 c9                xor    %r9d,%r9d
  400f91:       31 c9                   xor    %ecx,%ecx
  400f93:       48 8d 55 88             lea    -0x78(%rbp),%rdx
  400f97:       48 8b 75 e0             mov    -0x20(%rbp),%rsi
  400f9b:       48 89 70 28             mov    %rsi,0x28(%rax)
  400f9f:       48 8b 75 e8             mov    -0x18(%rbp),%rsi
  400fa3:       48 89 70 30             mov    %rsi,0x30(%rax)
  400fa7:       48 8b 75 d0             mov    -0x30(%rbp),%rsi
  400fab:       48 89 70 38             mov    %rsi,0x38(%rax)
  400faf:       48 8b 75 d8             mov    -0x28(%rbp),%rsi
  400fb3:       48 89 70 40             mov    %rsi,0x40(%rax)
  400fb7:       48 8b 34 25 40 1f 40    mov    0x401f40,%rsi
  400fbe:       00 
  400fbf:       48 89 70 48             mov    %rsi,0x48(%rax)
  400fc3:       48 8b 34 25 48 1f 40    mov    0x401f48,%rsi
  400fca:       00 
  400fcb:       48 89 70 50             mov    %rsi,0x50(%rax)
> 400fcf:       48 8b 31                mov    (%rcx),%rsi
  400fd2:       48 89 70 58             mov    %rsi,0x58(%rax)
  400fd6:       48 8b 49 08             mov    0x8(%rcx),%rcx
  400fda:       48 89 48 60             mov    %rcx,0x60(%rax)
  400fde:       48 b9 00 61 60 00 00    movabs $0x606100,%rcx
  400fe5:       00 00 00 
  400fe8:       48 89 4d 88             mov    %rcx,-0x78(%rbp)
  400fec:       48 c7 45 90 04 00 00    movq   $0x4,-0x70(%rbp)
  400ff3:       00 
  400ff4:       c6 45 98 03             movb   $0x3,-0x68(%rbp)
  400ff8:       48 c7 45 80 01 00 00    movq   $0x1,-0x80(%rbp)
  400fff:       00 
  401000:       48 8b 4d f8             mov    -0x8(%rbp),%rcx
  401004:       8b 31                   mov    (%rcx),%esi
  401006:       48 b9 97 1f 40 00 00    movabs $0x401f97,%rcx
  40100d:       00 00 00 
  401010:       48 89 4d b0             mov    %rcx,-0x50(%rbp)
  401014:       48 8d 7d a0             lea    -0x60(%rbp),%rdi
  401018:       48 89 95 18 fe ff ff    mov    %rdx,-0x1e8(%rbp)
  40101f:       48 89 c2                mov    %rax,%rdx
  401022:       b9 01 00 00 00          mov    $0x1,%ecx
  401027:       4c 8b 85 18 fe ff ff    mov    -0x1e8(%rbp),%r8
  40102e:       48 c7 04 24 00 00 00    movq   $0x0,(%rsp)
  401035:       00 
  401036:       e8 55 fb ff ff          callq  400b90 <__kmpc_omp_task_with_deps@plt>

[lildmh]

Can you check if https://reviews.llvm.org/D84470 fixes the problem as this test doesn't work on my machine? If yes, please accept and upstream it since I don't have the permission.

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;
 
+  /// Whether this clause is possible to have user-defined mappers associated.
+  /// It should be true for map, to, and from clauses, and false for
+  /// use_device_ptr and is_device_ptr.
+  bool hasMapper;
+
   /// C++ nested name specifier for the associated user-defined mapper.
   NestedNameSpecifierLoc MapperQualifierLoc;
 
   /// The associated user-defined mapper identifier information.
   DeclarationNameInfo MapperIdInfo;
 
 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 Locs Locations needed to build a mappable clause. It includes 1)
   /// StartLoc: starting location of the clause (the clause keyword); 2)
   /// LParenLoc: location of '('; 3) EndLoc: ending location of the clause.
   /// \param Sizes All required sizes to build a mappable clause. It includes 1)
   /// NumVars: number of expressions listed in this clause; 2)
   /// NumUniqueDeclarations: number of unique base declarations in this clause;
   /// 3) NumComponentLists: number of component lists in this clause; and 4)
   /// NumComponents: total number of expression components in the clause.
+  /// \param hasMapper Indicates whether this clause is possible to have
+  /// user-defined mappers associated.
   /// \param MapperQualifierLocPtr C++ nested name specifier for the associated
   /// user-defined mapper.
   /// \param MapperIdInfoPtr The identifier of associated user-defined mapper.
   OMPMappableExprListClause(
       OpenMPClauseKind K, const OMPVarListLocTy &Locs,
-      const OMPMappableExprListSizeTy &Sizes,
+      const OMPMappableExprListSizeTy &Sizes, bool hasMapper = false,
       NestedNameSpecifierLoc *MapperQualifierLocPtr = nullptr,
       DeclarationNameInfo *MapperIdInfoPtr = nullptr)
       : OMPVarListClause<T>(K, Locs.StartLoc, Locs.LParenLoc, Locs.EndLoc,
                             Sizes.NumVars),
         NumUniqueDeclarations(Sizes.NumUniqueDeclarations),
         NumComponentLists(Sizes.NumComponentLists),
-        NumComponents(Sizes.NumComponents) {
+        NumComponents(Sizes.NumComponents), hasMapper(hasMapper) {
     if (MapperQualifierLocPtr)
       MapperQualifierLoc = *MapperQualifierLocPtr;
     if (MapperIdInfoPtr)
       MapperIdInfo = *MapperIdInfoPtr;
   }
 
   /// Get the unique declarations that are in the trailing objects of the
   /// class.
protected:
   /// 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() {
+    assert(hasMapper &&
+           "Must be a clause that is possible to have user-defined mappers");
     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 *>() +
+  ArrayRef<const Expr *> getUDMapperRefs() const {
+    assert(hasMapper &&
+           "Must be a clause that is possible to have user-defined mappers");
+    return llvm::makeArrayRef<const Expr *>(
+        static_cast<const 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.");
+    assert(hasMapper &&
+           "Must be a clause that is possible to have 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.
class const_component_lists_iterator
             std::forward_iterator_tag, MappableComponent, ptrdiff_t,
             MappableComponent, MappableComponent> {
     // The declaration the iterator currently refers to.
     ArrayRef<ValueDecl *>::iterator DeclCur;
 
     // The list number associated with the current declaration.
     ArrayRef<unsigned>::iterator NumListsCur;
 
+    // Whether this clause is possible to have user-defined mappers associated.
+    bool hasMapper;
+
+    // The user-defined mapper associated with the current declaration.
+    ArrayRef<const Expr *>::iterator MapperCur;
+
     // Remaining lists for the current declaration.
     unsigned RemainingLists = 0;
 
     // The cumulative size of the previous list, or zero if there is no previous
     // list.
     unsigned PrevListSize = 0;
 
     // The cumulative sizes of the current list - it will delimit the remaining
     // range of interest.
     ArrayRef<unsigned>::const_iterator ListSizeCur;
     ArrayRef<unsigned>::const_iterator ListSizeEnd;
 
     // Iterator to the end of the components storage.
     MappableExprComponentListRef::const_iterator End;
 
   public:
     /// Construct an iterator that scans all lists.
     explicit const_component_lists_iterator(
         ArrayRef<ValueDecl *> UniqueDecls, ArrayRef<unsigned> DeclsListNum,
         ArrayRef<unsigned> CumulativeListSizes,
-        MappableExprComponentListRef Components)
+        MappableExprComponentListRef Components, bool hasMapper,
+        ArrayRef<const Expr *> Mappers)
         : const_component_lists_iterator::iterator_adaptor_base(
               Components.begin()),
           DeclCur(UniqueDecls.begin()), NumListsCur(DeclsListNum.begin()),
+          hasMapper(hasMapper), MapperCur(Mappers.begin()),
           ListSizeCur(CumulativeListSizes.begin()),
           ListSizeEnd(CumulativeListSizes.end()), End(Components.end()) {
       assert(UniqueDecls.size() == DeclsListNum.size() &&
              "Inconsistent number of declarations and list sizes!");
       if (!DeclsListNum.empty())
         RemainingLists = *NumListsCur;
     }
 
     /// Construct an iterator that scan lists for a given declaration \a
     /// Declaration.
     explicit const_component_lists_iterator(
         const ValueDecl *Declaration, ArrayRef<ValueDecl *> UniqueDecls,
         ArrayRef<unsigned> DeclsListNum, ArrayRef<unsigned> CumulativeListSizes,
-        MappableExprComponentListRef Components)
+        MappableExprComponentListRef Components, bool hasMapper,
+        ArrayRef<const Expr *> Mappers)
         : const_component_lists_iterator(UniqueDecls, DeclsListNum,
-                                         CumulativeListSizes, Components) {
+                                         CumulativeListSizes, Components,
+                                         hasMapper, Mappers) {
       // Look for the desired declaration. While we are looking for it, we
       // update the state so that we know the component where a given list
       // starts.
       for (; DeclCur != UniqueDecls.end(); ++DeclCur, ++NumListsCur) {
         if (*DeclCur == Declaration)
           break;
 
         assert(*NumListsCur > 0 && "No lists associated with declaration??");
 
         // Skip the lists associated with the current declaration, but save the
         // last list size that was skipped.
         std::advance(ListSizeCur, *NumListsCur - 1);
         PrevListSize = *ListSizeCur;
         ++ListSizeCur;
+
+        if (hasMapper)
+          ++MapperCur;
       }
 
       // If we didn't find any declaration, advance the iterator to after the
       // last component and set remaining lists to zero.
       if (ListSizeCur == CumulativeListSizes.end()) {
         this->I = End;
         RemainingLists = 0u;
         return;
explicit const_component_lists_iterator(
 
       // Given that the list sizes are cumulative, the index of the component
       // that start the list is the size of the previous list.
       std::advance(this->I, PrevListSize);
     }
 
     // Return the array with the current list. The sizes are cumulative, so the
     // array size is the difference between the current size and previous one.
-    std::pair<const ValueDecl *, MappableExprComponentListRef>
+    std::tuple<const ValueDecl *, MappableExprComponentListRef,
+               const ValueDecl *>
     operator*() const {
       assert(ListSizeCur != ListSizeEnd && "Invalid iterator!");
-      return std::make_pair(
+      const ValueDecl *Mapper = nullptr;
+      if (hasMapper && *MapperCur)
+        Mapper = cast<ValueDecl>(cast<DeclRefExpr>(*MapperCur)->getDecl());
+      return std::make_tuple(
           *DeclCur,
-          MappableExprComponentListRef(&*this->I, *ListSizeCur - PrevListSize));
+          MappableExprComponentListRef(&*this->I, *ListSizeCur - PrevListSize),
+          Mapper);
     }
-    std::pair<const ValueDecl *, MappableExprComponentListRef>
+    std::tuple<const ValueDecl *, MappableExprComponentListRef,
+               const ValueDecl *>
     operator->() const {
       return **this;
     }
 
     // Skip the components of the current list.
     const_component_lists_iterator &operator++() {
       assert(ListSizeCur != ListSizeEnd && RemainingLists &&
              "Invalid iterator!");
 
       // If we don't have more lists just skip all the components. Otherwise,
       // advance the iterator by the number of components in the current list.
       if (std::next(ListSizeCur) == ListSizeEnd) {
         this->I = End;
         RemainingLists = 0;
       } else {
         std::advance(this->I, *ListSizeCur - PrevListSize);
         PrevListSize = *ListSizeCur;
 
         // We are done with a declaration, move to the next one.
         if (!(--RemainingLists)) {
           ++DeclCur;
           ++NumListsCur;
+          if (hasMapper)
+            ++MapperCur;
           RemainingLists = *NumListsCur;
           assert(RemainingLists && "No lists in the following declaration??");
         }
       }
 
       ++ListSizeCur;
       return *this;
     }
   };
 
   using const_component_lists_range =
       llvm::iterator_range<const_component_lists_iterator>;
 
   /// Iterators for all component lists.
   const_component_lists_iterator component_lists_begin() const {
     return const_component_lists_iterator(
         getUniqueDeclsRef(), getDeclNumListsRef(), getComponentListSizesRef(),
-        getComponentsRef());
+        getComponentsRef(), hasMapper,
+        hasMapper ? getUDMapperRefs() : ArrayRef<const Expr *>());
   }
   const_component_lists_iterator component_lists_end() const {
     return const_component_lists_iterator(
         ArrayRef<ValueDecl *>(), ArrayRef<unsigned>(), ArrayRef<unsigned>(),
         MappableExprComponentListRef(getComponentsRef().end(),
-                                     getComponentsRef().end()));
+                                     getComponentsRef().end()),
+        hasMapper, ArrayRef<const Expr *>());
   }
   const_component_lists_range component_lists() const {
     return {component_lists_begin(), component_lists_end()};
   }
 
   /// Iterators for component lists associated with the provided
   /// declaration.
   const_component_lists_iterator
   decl_component_lists_begin(const ValueDecl *VD) const {
     return const_component_lists_iterator(
         VD, getUniqueDeclsRef(), getDeclNumListsRef(),
-        getComponentListSizesRef(), getComponentsRef());
+        getComponentListSizesRef(), getComponentsRef(), hasMapper,
+        hasMapper ? getUDMapperRefs() : ArrayRef<const Expr *>());
   }
   const_component_lists_iterator decl_component_lists_end() const {
     return component_lists_end();
   }
   const_component_lists_range decl_component_lists(const ValueDecl *VD) const {
     return {decl_component_lists_begin(VD), decl_component_lists_end()};
   }
 
private:
   /// 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, const OMPVarListLocTy &Locs,
                         const OMPMappableExprListSizeTy &Sizes)
-      : OMPMappableExprListClause(OMPC_map, Locs, Sizes, &MapperQualifierLoc,
-                                  &MapperIdInfo),
+      : OMPMappableExprListClause(OMPC_map, Locs, Sizes, /*hasMapper=*/true,
+                                  &MapperQualifierLoc, &MapperIdInfo),
         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));
 
     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 Sizes All required sizes to build a mappable clause. It includes 1)
   /// NumVars: number of expressions listed in this clause; 2)
   /// NumUniqueDeclarations: number of unique base declarations in this clause;
   /// 3) NumComponentLists: number of component lists in this clause; and 4)
   /// NumComponents: total number of expression components in the clause.
   explicit OMPMapClause(const OMPMappableExprListSizeTy &Sizes)
-      : OMPMappableExprListClause(OMPC_map, OMPVarListLocTy(), Sizes) {}
+      : OMPMappableExprListClause(OMPC_map, OMPVarListLocTy(), Sizes,
+                                  /*hasMapper=*/true) {}
 
   /// 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.");
class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
   /// 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 OMPToClause(NestedNameSpecifierLoc MapperQualifierLoc,
                        DeclarationNameInfo MapperIdInfo,
                        const OMPVarListLocTy &Locs,
                        const OMPMappableExprListSizeTy &Sizes)
-      : OMPMappableExprListClause(OMPC_to, Locs, Sizes, &MapperQualifierLoc,
-                                  &MapperIdInfo) {}
+      : OMPMappableExprListClause(OMPC_to, Locs, Sizes, /*hasMapper=*/true,
+                                  &MapperQualifierLoc, &MapperIdInfo) {}
 
   /// Build an empty clause.
   ///
   /// \param Sizes All required sizes to build a mappable clause. It includes 1)
   /// NumVars: number of expressions listed in this clause; 2)
   /// NumUniqueDeclarations: number of unique base declarations in this clause;
   /// 3) NumComponentLists: number of component lists in this clause; and 4)
   /// NumComponents: total number of expression components in the clause.
   explicit OMPToClause(const OMPMappableExprListSizeTy &Sizes)
-      : OMPMappableExprListClause(OMPC_to, OMPVarListLocTy(), Sizes) {}
+      : OMPMappableExprListClause(OMPC_to, OMPVarListLocTy(), Sizes,
+                                  /*hasMapper=*/true) {}
 
   /// 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 {
     // There are varlist_size() of expressions, and varlist_size() of
     // user-defined mappers.
     return 2 * varlist_size();
   }
class OMPFromClause final
   /// NumVars: number of expressions listed in this clause; 2)
   /// NumUniqueDeclarations: number of unique base declarations in this clause;
   /// 3) NumComponentLists: number of component lists in this clause; and 4)
   /// NumComponents: total number of expression components in the clause.
   explicit OMPFromClause(NestedNameSpecifierLoc MapperQualifierLoc,
                          DeclarationNameInfo MapperIdInfo,
                          const OMPVarListLocTy &Locs,
                          const OMPMappableExprListSizeTy &Sizes)
-      : OMPMappableExprListClause(OMPC_from, Locs, Sizes, &MapperQualifierLoc,
-                                  &MapperIdInfo) {}
+      : OMPMappableExprListClause(OMPC_from, Locs, Sizes, /*hasMapper=*/true,
+                                  &MapperQualifierLoc, &MapperIdInfo) {}
 
   /// Build an empty clause.
   ///
   /// \param Sizes All required sizes to build a mappable clause. It includes 1)
   /// NumVars: number of expressions listed in this clause; 2)
   /// NumUniqueDeclarations: number of unique base declarations in this clause;
   /// 3) NumComponentLists: number of component lists in this clause; and 4)
   /// NumComponents: total number of expression components in the clause.
   explicit OMPFromClause(const OMPMappableExprListSizeTy &Sizes)
-      : OMPMappableExprListClause(OMPC_from, OMPVarListLocTy(), Sizes) {}
+      : OMPMappableExprListClause(OMPC_from, OMPVarListLocTy(), Sizes,
+                                  /*hasMapper=*/true) {}
 
   /// 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 {
     // There are varlist_size() of expressions, and varlist_size() of
     // user-defined mappers.
     return 2 * varlist_size();
   }

lib/CodeGen/CGOpenMPRuntime.h

virtual void emitUserDefinedReduction(CodeGenFunction *CGF,
                                         const OMPDeclareReductionDecl *D);
   /// Get combiner/initializer for the specified user-defined reduction, if any.
   virtual std::pair<llvm::Function *, llvm::Function *>
   getUserDefinedReduction(const OMPDeclareReductionDecl *D);
 
   /// Emit the function for the user defined mapper construct.
   void emitUserDefinedMapper(const OMPDeclareMapperDecl *D,
                              CodeGenFunction *CGF = nullptr);
+  /// Get the function for the specified user-defined mapper, if any.
+  virtual llvm::Function *
+  getUserDefinedMapperFunc(const OMPDeclareMapperDecl *D);
 
   /// Emits outlined function for the specified OpenMP parallel directive
   /// \a D. This outlined function has type void(*)(kmp_int32 *ThreadID,
   /// kmp_int32 BoundID, struct context_vars*).
   /// \param D OpenMP directive.
   /// \param ThreadIDVar Variable for thread id in the current OpenMP region.
   /// \param InnermostKind Kind of innermost directive (for simple directives it
   /// is a directive itself, for combined - its innermost directive).
public:
     /// The array of base pointer passed to the runtime library.
     llvm::Value *BasePointersArray = nullptr;
     /// The array of section pointers passed to the runtime library.
     llvm::Value *PointersArray = nullptr;
     /// The array of sizes passed to the runtime library.
     llvm::Value *SizesArray = nullptr;
     /// The array of map types passed to the runtime library.
     llvm::Value *MapTypesArray = nullptr;
+    /// The array of user-defined mappers passed to the runtime library.
+    llvm::Value *MappersArray = nullptr;
     /// The total number of pointers passed to the runtime library.
     unsigned NumberOfPtrs = 0u;
     /// Map between the a declaration of a capture and the corresponding base
     /// pointer address where the runtime returns the device pointers.
     llvm::DenseMap<const ValueDecl *, Address> CaptureDeviceAddrMap;
 
     explicit TargetDataInfo() {}
     explicit TargetDataInfo(bool RequiresDevicePointerInfo)
         : RequiresDevicePointerInfo(RequiresDevicePointerInfo) {}
     /// Clear information about the data arrays.
     void clearArrayInfo() {
       BasePointersArray = nullptr;
       PointersArray = nullptr;
       SizesArray = nullptr;
       MapTypesArray = nullptr;
+      MappersArray = nullptr;
       NumberOfPtrs = 0u;
     }
     /// Return true if the current target data information has valid arrays.
     bool isValid() {
       return BasePointersArray && PointersArray && SizesArray &&
-             MapTypesArray && NumberOfPtrs;
+             MapTypesArray && MappersArray && NumberOfPtrs;
     }
     bool requiresDevicePointerInfo() { return RequiresDevicePointerInfo; }
   };
 
   /// Emit the target data mapping code associated with \a D.
   /// \param D Directive to emit.
   /// \param IfCond Expression evaluated in if clause associated with the
   /// target directive, or null if no device clause is used.

lib/CodeGen/CGOpenMPRuntime.cpp

 #include "llvm/ADT/ArrayRef.h"
 #include "llvm/Bitcode/BitcodeReader.h"
 #include "llvm/IR/DerivedTypes.h"
 #include "llvm/IR/GlobalValue.h"
 #include "llvm/IR/Value.h"
 #include "llvm/Support/Format.h"
 #include "llvm/Support/raw_ostream.h"
 #include <cassert>
+#include <iostream>
 
 using namespace clang;
 using namespace CodeGen;
 
 namespace {
 /// Base class for handling code generation inside OpenMP regions.
 class CGOpenMPRegionInfo : public CodeGenFunction::CGCapturedStmtInfo {
 public:
enum OpenMPRTLFunction {
   // Call to void __tgt_target_data_begin(int64_t device_id, int32_t arg_num,
   // void** args_base, void **args, int64_t *arg_sizes, int64_t *arg_types);
   OMPRTL__tgt_target_data_begin,
   // Call to void __tgt_target_data_begin_nowait(int64_t device_id, int32_t
   // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
   // *arg_types);
   OMPRTL__tgt_target_data_begin_nowait,
   // Call to void __tgt_target_data_end(int64_t device_id, int32_t arg_num,
-  // void** args_base, void **args, size_t *arg_sizes, int64_t *arg_types);
+  // void** args_base, void **args, int64_t *arg_sizes, int64_t *arg_types);
   OMPRTL__tgt_target_data_end,
   // Call to void __tgt_target_data_end_nowait(int64_t device_id, int32_t
   // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
   // *arg_types);
   OMPRTL__tgt_target_data_end_nowait,
   // Call to void __tgt_target_data_update(int64_t device_id, int32_t arg_num,
   // void** args_base, void **args, int64_t *arg_sizes, int64_t *arg_types);
   OMPRTL__tgt_target_data_update,
   // Call to void __tgt_target_data_update_nowait(int64_t device_id, int32_t
   // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
   // *arg_types);
   OMPRTL__tgt_target_data_update_nowait,
+  // Call to int32_t __tgt_target_mapper(int64_t device_id, void *host_ptr,
+  // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+  // *arg_types, void **arg_mappers);
+  OMPRTL__tgt_target_mapper,
+  // Call to int32_t __tgt_target_nowait_mapper(int64_t device_id, void
+  // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t
+  // *arg_sizes, int64_t *arg_types, void **arg_mappers);
+  OMPRTL__tgt_target_nowait_mapper,
+  // Call to int32_t __tgt_target_teams_mapper(int64_t device_id, void
+  // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t
+  // *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t num_teams,
+  // int32_t thread_limit);
+  OMPRTL__tgt_target_teams_mapper,
+  // Call to int32_t __tgt_target_teams_nowait_mapper(int64_t device_id, void
+  // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t
+  // *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t num_teams,
+  // int32_t thread_limit);
+  OMPRTL__tgt_target_teams_nowait_mapper,
+  // Call to void __tgt_target_data_begin_mapper(int64_t device_id, int32_t
+  // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+  // *arg_types, void **arg_mappers);
+  OMPRTL__tgt_target_data_begin_mapper,
+  // Call to void __tgt_target_data_begin_nowait_mapper(int64_t device_id,
+  // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+  // *arg_types, void **arg_mappers);
+  OMPRTL__tgt_target_data_begin_nowait_mapper,
+  // Call to void __tgt_target_data_end_mapper(int64_t device_id, int32_t
+  // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+  // *arg_types, void **arg_mappers);
+  OMPRTL__tgt_target_data_end_mapper,
+  // Call to void __tgt_target_data_end_nowait_mapper(int64_t device_id, int32_t
+  // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+  // *arg_types, void **arg_mappers);
+  OMPRTL__tgt_target_data_end_nowait_mapper,
+  // Call to void __tgt_target_data_update_mapper(int64_t device_id, int32_t
+  // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+  // *arg_types, void **arg_mappers);
+  OMPRTL__tgt_target_data_update_mapper,
+  // Call to void __tgt_target_data_update_nowait_mapper(int64_t device_id,
+  // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+  // *arg_types, void **arg_mappers);
+  OMPRTL__tgt_target_data_update_nowait_mapper,
   // Call to int64_t __tgt_mapper_num_components(void *rt_mapper_handle);
   OMPRTL__tgt_mapper_num_components,
   // Call to void __tgt_push_mapper_component(void *rt_mapper_handle, void
   // *base, void *begin, int64_t size, int64_t type);
   OMPRTL__tgt_push_mapper_component,
 };
 
 /// A basic class for pre|post-action for advanced codegen sequence for OpenMP
llvm::Type *TypeParams[] = {CGM.Int64Ty,
                                 CGM.VoidPtrPtrTy,
                                 CGM.Int64Ty->getPointerTo(),
                                 CGM.Int64Ty->getPointerTo()};
     auto *FnTy =
         llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false);
     RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_update_nowait");
     break;
   }
+  case OMPRTL__tgt_target_mapper: {
+    // Build int32_t __tgt_target_mapper(int64_t device_id, void *host_ptr,
+    // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes,
+    // int64_t *arg_types, void **arg_mappers);
+    llvm::Type *TypeParams[] = {CGM.Int64Ty,
+                                CGM.VoidPtrTy,
+                                CGM.Int32Ty,
+                                CGM.VoidPtrPtrTy,
+                                CGM.VoidPtrPtrTy,
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.VoidPtrPtrTy};
+    auto *FnTy =
+        llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_mapper");
+    break;
+  }
+  case OMPRTL__tgt_target_nowait_mapper: {
+    // Build int32_t __tgt_target_nowait_mapper(int64_t device_id, void
+    // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t
+    // *arg_sizes, int64_t *arg_types, void **arg_mappers);
+    llvm::Type *TypeParams[] = {CGM.Int64Ty,
+                                CGM.VoidPtrTy,
+                                CGM.Int32Ty,
+                                CGM.VoidPtrPtrTy,
+                                CGM.VoidPtrPtrTy,
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.VoidPtrPtrTy};
+    auto *FnTy =
+        llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_nowait_mapper");
+    break;
+  }
+  case OMPRTL__tgt_target_teams_mapper: {
+    // Build int32_t __tgt_target_teams_mapper(int64_t device_id, void
+    // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t
+    // *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t num_teams,
+    // int32_t thread_limit);
+    llvm::Type *TypeParams[] = {CGM.Int64Ty,
+                                CGM.VoidPtrTy,
+                                CGM.Int32Ty,
+                                CGM.VoidPtrPtrTy,
+                                CGM.VoidPtrPtrTy,
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.VoidPtrPtrTy,
+                                CGM.Int32Ty,
+                                CGM.Int32Ty};
+    auto *FnTy =
+        llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_teams_mapper");
+    break;
+  }
+  case OMPRTL__tgt_target_teams_nowait_mapper: {
+    // Build int32_t __tgt_target_teams_nowait_mapper(int64_t device_id, void
+    // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t
+    // *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t num_teams,
+    // int32_t thread_limit);
+    llvm::Type *TypeParams[] = {CGM.Int64Ty,
+                                CGM.VoidPtrTy,
+                                CGM.Int32Ty,
+                                CGM.VoidPtrPtrTy,
+                                CGM.VoidPtrPtrTy,
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.VoidPtrPtrTy,
+                                CGM.Int32Ty,
+                                CGM.Int32Ty};
+    auto *FnTy =
+        llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_teams_nowait_mapper");
+    break;
+  }
+  case OMPRTL__tgt_target_data_begin_mapper: {
+    // Build void __tgt_target_data_begin_mapper(int64_t device_id, int32_t
+    // arg_num, void **args_base, void **args, int64_t *arg_sizes, int64_t
+    // *arg_types, void **arg_mappers);
+    llvm::Type *TypeParams[] = {CGM.Int64Ty,
+                                CGM.Int32Ty,
+                                CGM.VoidPtrPtrTy,
+                                CGM.VoidPtrPtrTy,
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.VoidPtrPtrTy};
+    auto *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_begin_mapper");
+    break;
+  }
+  case OMPRTL__tgt_target_data_begin_nowait_mapper: {
+    // Build void __tgt_target_data_begin_nowait_mapper(int64_t device_id,
+    // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes,
+    // int64_t *arg_types, void **arg_mappers);
+    llvm::Type *TypeParams[] = {CGM.Int64Ty,
+                                CGM.Int32Ty,
+                                CGM.VoidPtrPtrTy,
+                                CGM.VoidPtrPtrTy,
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.VoidPtrPtrTy};
+    auto *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy,
+                                      "__tgt_target_data_begin_nowait_mapper");
+    break;
+  }
+  case OMPRTL__tgt_target_data_end_mapper: {
+    // Build void __tgt_target_data_end_mapper(int64_t device_id, int32_t
+    // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+    // *arg_types, void **arg_mappers);
+    llvm::Type *TypeParams[] = {CGM.Int64Ty,
+                                CGM.Int32Ty,
+                                CGM.VoidPtrPtrTy,
+                                CGM.VoidPtrPtrTy,
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.VoidPtrPtrTy};
+    auto *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_end_mapper");
+    break;
+  }
+  case OMPRTL__tgt_target_data_end_nowait_mapper: {
+    // Build void __tgt_target_data_end_nowait_mapper(int64_t device_id, int32_t
+    // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+    // *arg_types, void **arg_mappers);
+    llvm::Type *TypeParams[] = {CGM.Int64Ty,
+                                CGM.Int32Ty,
+                                CGM.VoidPtrPtrTy,
+                                CGM.VoidPtrPtrTy,
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.VoidPtrPtrTy};
+    auto *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false);
+    RTLFn =
+        CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_end_nowait_mapper");
+    break;
+  }
+  case OMPRTL__tgt_target_data_update_mapper: {
+    // Build void __tgt_target_data_update_mapper(int64_t device_id, int32_t
+    // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t
+    // *arg_types, void **arg_mappers);
+    llvm::Type *TypeParams[] = {CGM.Int64Ty,
+                                CGM.Int32Ty,
+                                CGM.VoidPtrPtrTy,
+                                CGM.VoidPtrPtrTy,
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.VoidPtrPtrTy};
+    auto *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_update_mapper");
+    break;
+  }
+  case OMPRTL__tgt_target_data_update_nowait_mapper: {
+    // Build void __tgt_target_data_update_nowait_mapper(int64_t device_id,
+    // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes,
+    // int64_t *arg_types, void **arg_mappers);
+    llvm::Type *TypeParams[] = {CGM.Int64Ty,
+                                CGM.Int32Ty,
+                                CGM.VoidPtrPtrTy,
+                                CGM.VoidPtrPtrTy,
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.VoidPtrPtrTy};
+    auto *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy,
+                                      "__tgt_target_data_update_nowait_mapper");
+    break;
+  }
   case OMPRTL__tgt_mapper_num_components: {
     // Build int64_t __tgt_mapper_num_components(void *rt_mapper_handle);
     llvm::Type *TypeParams[] = {CGM.VoidPtrTy};
     auto *FnTy =
         llvm::FunctionType::get(CGM.Int64Ty, TypeParams, /*isVarArg*/ false);
     RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_mapper_num_components");
     break;
   }
public:
     llvm::Value *operator*() const { return Ptr; }
     const ValueDecl *getDevicePtrDecl() const { return DevPtrDecl; }
     void setDevicePtrDecl(const ValueDecl *D) { DevPtrDecl = D; }
   };
 
   using MapBaseValuesArrayTy = SmallVector<BasePointerInfo, 4>;
   using MapValuesArrayTy = SmallVector<llvm::Value *, 4>;
   using MapFlagsArrayTy = SmallVector<OpenMPOffloadMappingFlags, 4>;
+  using MapMappersArrayTy = SmallVector<const ValueDecl *, 4>;
 
   /// Map between a struct and the its lowest & highest elements which have been
   /// mapped.
   /// [ValueDecl *] --> {LE(FieldIndex, Pointer),
   ///                    HE(FieldIndex, Pointer)}
   struct StructRangeInfoTy {
     std::pair<unsigned /*FieldIndex*/, Address /*Pointer*/> LowestElem = {
         0, Address::invalid()};
     std::pair<unsigned /*FieldIndex*/, Address /*Pointer*/> HighestElem = {
         0, Address::invalid()};
     Address Base = Address::invalid();
   };
 
 private:
   /// Kind that defines how a device pointer has to be returned.
   struct MapInfo {
     OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
     OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
     ArrayRef<OpenMPMapModifierKind> MapModifiers;
     bool ReturnDevicePointer = false;
     bool IsImplicit = false;
+    const ValueDecl *Mapper = nullptr;
 
     MapInfo() = default;
     MapInfo(
         OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
         OpenMPMapClauseKind MapType,
-        ArrayRef<OpenMPMapModifierKind> MapModifiers,
-        bool ReturnDevicePointer, bool IsImplicit)
+        ArrayRef<OpenMPMapModifierKind> MapModifiers, bool ReturnDevicePointer,
+        bool IsImplicit, const ValueDecl *Mapper = nullptr)
         : Components(Components), MapType(MapType), MapModifiers(MapModifiers),
-          ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit) {}
+          ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit),
+          Mapper(Mapper) {}
   };
 
   /// If use_device_ptr is used on a pointer which is a struct member and there
   /// is no map information about it, then emission of that entry is deferred
   /// until the whole struct has been processed.
   struct DeferredDevicePtrEntryTy {
     const Expr *IE = nullptr;
     const ValueDecl *VD = nullptr;
private:
 
   /// Map between device pointer declarations and their expression components.
   /// The key value for declarations in 'this' is null.
   llvm::DenseMap<
       const ValueDecl *,
       SmallVector<OMPClauseMappableExprCommon::MappableExprComponentListRef, 4>>
       DevPointersMap;
 
-  llvm::Value *getExprTypeSize(const Expr *E) const {
+  llvm::Value *getExprTypeSize(const Expr *E, bool hasMapper) const {
     QualType ExprTy = E->getType().getCanonicalType();
 
     // Reference types are ignored for mapping purposes.
     if (const auto *RefTy = ExprTy->getAs<ReferenceType>())
       ExprTy = RefTy->getPointeeType().getCanonicalType();
 
     // Given that an array section is considered a built-in type, we need to
     // do the calculation based on the length of the section instead of relying
     // on CGF.getTypeSize(E->getType()).
     if (const auto *OAE = dyn_cast<OMPArraySectionExpr>(E)) {
       QualType BaseTy = OMPArraySectionExpr::getBaseOriginalType(
                             OAE->getBase()->IgnoreParenImpCasts())
                             .getCanonicalType();
 
       // If there is no length associated with the expression, that means we
       // are using the whole length of the base.
-      if (!OAE->getLength() && OAE->getColonLoc().isValid())
+      if (!OAE->getLength() && OAE->getColonLoc().isValid()) {
+        // In case that a user-defined mapper is attached, its size is the
+        // number of array elements instead of the number of total bytes.
+        if (hasMapper)
+          return CGF.Builder.getInt64(1);
+        else
           return CGF.getTypeSize(BaseTy);
+      }
 
       llvm::Value *ElemSize;
       if (const auto *PTy = BaseTy->getAs<PointerType>()) {
         ElemSize = CGF.getTypeSize(PTy->getPointeeType().getCanonicalType());
       } else {
         const auto *ATy = cast<ArrayType>(BaseTy.getTypePtr());
         assert(ATy && "Expecting array type if not a pointer type.");
         ElemSize = CGF.getTypeSize(ATy->getElementType().getCanonicalType());
       }
 
       // If we don't have a length at this point, that is because we have an
       // array section with a single element.
-      if (!OAE->getLength())
+      if (!OAE->getLength()) {
+        // In case that a user-defined mapper is attached, its size is the
+        // number of array elements instead of the number of total bytes.
+        if (hasMapper)
+          return CGF.Builder.getInt64(1);
+        else
           return ElemSize;
+      }
 
       llvm::Value *LengthVal = CGF.EmitScalarExpr(OAE->getLength());
       LengthVal =
           CGF.Builder.CreateIntCast(LengthVal, CGF.SizeTy, /*isSigned=*/false);
+      // In case that a user-defined mapper is attached, its size is the
+      // number of array elements instead of the number of total bytes.
+      if (hasMapper)
+        return LengthVal;
+      else
         return CGF.Builder.CreateNUWMul(LengthVal, ElemSize);
     }
+    // In case that a user-defined mapper is attached, its size is the
+    // number of array elements instead of the number of total bytes.
+    if (hasMapper)
+      return CGF.Builder.getInt64(1);
+    else
       return CGF.getTypeSize(ExprTy);
   }
 
   /// Return the corresponding bits for a given map clause modifier. Add
   /// a flag marking the map as a pointer if requested. Add a flag marking the
   /// map as the first one of a series of maps that relate to the same map
   /// expression.
   OpenMPOffloadMappingFlags getMapTypeBits(
       OpenMPMapClauseKind MapType, ArrayRef<OpenMPMapModifierKind> MapModifiers,
bool isFinalArraySectionExpression(const Expr *E) const {
     Expr::EvalResult Result;
     if (!Length->EvaluateAsInt(Result, CGF.getContext()))
       return true; // Can have more that size 1.
 
     llvm::APSInt ConstLength = Result.Val.getInt();
     return ConstLength.getSExtValue() != 1;
   }
 
-  /// Generate the base pointers, section pointers, sizes and map type
-  /// bits for the provided map type, map modifier, and expression components.
-  /// \a IsFirstComponent should be set to true if the provided set of
-  /// components is the first associated with a capture.
+  /// Generate the base pointers, section pointers, sizes, map type bits, and
+  /// mappers for the provided map type, map modifier, and expression
+  /// components. \a IsFirstComponent should be set to true if the provided set
+  /// of components is the first associated with a capture.
   void generateInfoForComponentList(
-      OpenMPMapClauseKind MapType,
-      ArrayRef<OpenMPMapModifierKind> MapModifiers,
+      OpenMPMapClauseKind MapType, ArrayRef<OpenMPMapModifierKind> MapModifiers,
       OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
       MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers,
       MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
-      StructRangeInfoTy &PartialStruct, bool IsFirstComponentList,
-      bool IsImplicit,
+      MapMappersArrayTy &Mappers, StructRangeInfoTy &PartialStruct,
+      bool IsFirstComponentList, bool IsImplicit,
+      const ValueDecl *Mapper = nullptr,
       ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
           OverlappedElements = llvm::None) const {
     // The following summarizes what has to be generated for each map and the
     // types below. The generated information is expressed in this order:
     // base pointer, section pointer, size, flags
     // (to add to the ones that come from the map type and modifier).
     //
     // double d;
     // int i[100];
     // float *p;
for (; I != CE; ++I) {
                 break;
               }
             }
             BasePointers.push_back(BP.getPointer());
             Pointers.push_back(LB.getPointer());
             Sizes.push_back(CGF.Builder.CreateIntCast(Size, CGF.Int64Ty,
                                                       /*isSigned=*/true));
             Types.push_back(Flags);
+            Mappers.push_back(nullptr);
             LB = CGF.Builder.CreateConstGEP(ComponentLB, 1);
           }
           BasePointers.push_back(BP.getPointer());
           Pointers.push_back(LB.getPointer());
           Size = CGF.Builder.CreatePtrDiff(
               CGF.EmitCastToVoidPtr(
                   CGF.Builder.CreateConstGEP(HB, 1).getPointer()),
               CGF.EmitCastToVoidPtr(LB.getPointer()));
           Sizes.push_back(
               CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
           Types.push_back(Flags);
+          Mappers.push_back(nullptr);
           break;
         }
-        llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression());
         if (!IsMemberPointer) {
           BasePointers.push_back(BP.getPointer());
           Pointers.push_back(LB.getPointer());
-          Sizes.push_back(
-              CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
+
+          // If Mapper is valid, the last component inherits the mapper.
+          bool hasMapper = Mapper && Next == CE;
+          llvm::Value *Size =
+              getExprTypeSize(I->getAssociatedExpression(), hasMapper);
+          Sizes.push_back(CGF.Builder.CreateIntCast(Size, CGF.Int64Ty,
+                                                    /*isSigned=*/true));
+          if (hasMapper)
+            Mappers.push_back(Mapper);
+          else
+            Mappers.push_back(nullptr);
+          std::cerr << "HH " << "\n";
+          if (hasMapper) {
+            std::cerr << "AM: " << Mappers.size() << " ";
+            I->getAssociatedExpression()->dump();
+            //Size->dump();
+          }
 
           // We need to add a pointer flag for each map that comes from the
           // same expression except for the first one. We also need to signal
           // this map is the first one that relates with the current capture
           // (there is a set of entries for each capture).
           OpenMPOffloadMappingFlags Flags = getMapTypeBits(
               MapType, MapModifiers, IsImplicit,
               !IsExpressionFirstInfo || RequiresReference,
MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF)
     // Extract firstprivate clause information.
     for (const auto *C : Dir.getClausesOfKind<OMPFirstprivateClause>())
       for (const auto *D : C->varlists())
         FirstPrivateDecls.try_emplace(
             cast<VarDecl>(cast<DeclRefExpr>(D)->getDecl()), C->isImplicit());
     // Extract device pointer clause information.
     for (const auto *C : Dir.getClausesOfKind<OMPIsDevicePtrClause>())
       for (auto L : C->component_lists())
-        DevPointersMap[L.first].push_back(L.second);
+        DevPointersMap[std::get<0>(L)].push_back(std::get<1>(L));
   }
 
   /// Constructor for the declare mapper directive.
   MappableExprsHandler(const OMPDeclareMapperDecl &Dir, CodeGenFunction &CGF)
       : CurDir(&Dir), CGF(CGF) {}
 
   /// Generate code for the combined entry if we have a partially mapped struct
   /// and take care of the mapping flags of the arguments corresponding to
   /// individual struct members.
   void emitCombinedEntry(MapBaseValuesArrayTy &BasePointers,
                          MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes,
-                         MapFlagsArrayTy &Types, MapFlagsArrayTy &CurTypes,
+                         MapFlagsArrayTy &Types, MapMappersArrayTy &Mappers,
+                         MapFlagsArrayTy &CurTypes,
                          const StructRangeInfoTy &PartialStruct) const {
     // Base is the base of the struct
     BasePointers.push_back(PartialStruct.Base.getPointer());
     // Pointer is the address of the lowest element
     llvm::Value *LB = PartialStruct.LowestElem.second.getPointer();
     Pointers.push_back(LB);
+    // There should not be a mapper for a combined entry.
+    Mappers.push_back(nullptr);
     // Size is (addr of {highest+1} element) - (addr of lowest element)
     llvm::Value *HB = PartialStruct.HighestElem.second.getPointer();
     llvm::Value *HAddr = CGF.Builder.CreateConstGEP1_32(HB, /*Idx0=*/1);
     llvm::Value *CLAddr = CGF.Builder.CreatePointerCast(LB, CGF.VoidPtrTy);
     llvm::Value *CHAddr = CGF.Builder.CreatePointerCast(HAddr, CGF.VoidPtrTy);
     llvm::Value *Diff = CGF.Builder.CreatePtrDiff(CHAddr, CLAddr);
     llvm::Value *Size = CGF.Builder.CreateIntCast(Diff, CGF.Int64Ty,
                                                   /*isSigned=*/false);
     Sizes.push_back(Size);
     // Map type is always TARGET_PARAM
     Types.push_back(OMP_MAP_TARGET_PARAM);
     // Remove TARGET_PARAM flag from the first element
     (*CurTypes.begin()) &= ~OMP_MAP_TARGET_PARAM;
 
     // All other current entries will be MEMBER_OF the combined entry
     // (except for PTR_AND_OBJ entries which do not have a placeholder value
     // 0xFFFF in the MEMBER_OF field).
     OpenMPOffloadMappingFlags MemberOfFlag =
         getMemberOfFlag(BasePointers.size() - 1);
     for (auto &M : CurTypes)
       setCorrectMemberOfFlag(M, MemberOfFlag);
   }
 
-  /// Generate all the base pointers, section pointers, sizes and map
-  /// types for the extracted mappable expressions. Also, for each item that
+  /// Generate all the base pointers, section pointers, sizes, map types, and
+  /// mappers for the extracted mappable expressions. Also, for each item that
   /// relates with a device pointer, a pair of the relevant declaration and
   /// index where it occurs is appended to the device pointers info array.
   void generateAllInfo(MapBaseValuesArrayTy &BasePointers,
                        MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes,
-                       MapFlagsArrayTy &Types) const {
+                       MapFlagsArrayTy &Types,
+                       MapMappersArrayTy &Mappers) const {
     // We have to process the component lists that relate with the same
     // declaration in a single chunk so that we can generate the map flags
     // correctly. Therefore, we organize all lists in a map.
     llvm::MapVector<const ValueDecl *, SmallVector<MapInfo, 8>> Info;
 
     // Helper function to fill the information map for the different supported
     // clauses.
-    auto &&InfoGen = [&Info](
-        const ValueDecl *D,
+    auto &&InfoGen =
+        [&Info](const ValueDecl *D,
                 OMPClauseMappableExprCommon::MappableExprComponentListRef L,
                 OpenMPMapClauseKind MapType,
                 ArrayRef<OpenMPMapModifierKind> MapModifiers,
-        bool ReturnDevicePointer, bool IsImplicit) {
+                bool ReturnDevicePointer, bool IsImplicit,
+                const ValueDecl *Mapper) {
           const ValueDecl *VD =
               D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
           Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer,
-                            IsImplicit);
+                                IsImplicit, Mapper);
         };
 
     assert(CurDir.is<const OMPExecutableDirective *>() &&
            "Expect a executable directive");
     const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
     for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>())
       for (const auto &L : C->component_lists()) {
-        InfoGen(L.first, L.second, C->getMapType(), C->getMapTypeModifiers(),
-            /*ReturnDevicePointer=*/false, C->isImplicit());
+        InfoGen(std::get<0>(L), std::get<1>(L), C->getMapType(),
+                C->getMapTypeModifiers(), /*ReturnDevicePointer=*/false,
+                C->isImplicit(), std::get<2>(L));
       }
     for (const auto *C : CurExecDir->getClausesOfKind<OMPToClause>())
       for (const auto &L : C->component_lists()) {
-        InfoGen(L.first, L.second, OMPC_MAP_to, llvm::None,
-            /*ReturnDevicePointer=*/false, C->isImplicit());
+        InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_to, llvm::None,
+                /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L));
       }
     for (const auto *C : CurExecDir->getClausesOfKind<OMPFromClause>())
       for (const auto &L : C->component_lists()) {
-        InfoGen(L.first, L.second, OMPC_MAP_from, llvm::None,
-            /*ReturnDevicePointer=*/false, C->isImplicit());
+        InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_from, llvm::None,
+                /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L));
       }
 
     // Look at the use_device_ptr clause information and mark the existing map
     // entries as such. If there is no map information for an entry in the
     // use_device_ptr list, we create one with map type 'alloc' and zero size
     // section. It is the user fault if that was not mapped before. If there is
     // no map information and the pointer is a struct member, then we defer the
     // emission of that entry until the whole struct has been processed.
     llvm::MapVector<const ValueDecl *, SmallVector<DeferredDevicePtrEntryTy, 4>>
         DeferredInfo;
 
     for (const auto *C :
          CurExecDir->getClausesOfKind<OMPUseDevicePtrClause>()) {
       for (const auto &L : C->component_lists()) {
-        assert(!L.second.empty() && "Not expecting empty list of components!");
-        const ValueDecl *VD = L.second.back().getAssociatedDeclaration();
+        OMPClauseMappableExprCommon::MappableExprComponentListRef Components =
+            std::get<1>(L);
+        assert(!Components.empty() &&
+               "Not expecting empty list of components!");
+        const ValueDecl *VD = Components.back().getAssociatedDeclaration();
         VD = cast<ValueDecl>(VD->getCanonicalDecl());
-        const Expr *IE = L.second.back().getAssociatedExpression();
+        const Expr *IE = Components.back().getAssociatedExpression();
         // If the first component is a member expression, we have to look into
         // 'this', which maps to null in the map of map information. Otherwise
         // look directly for the information.
         auto It = Info.find(isa<MemberExpr>(IE) ? nullptr : VD);
 
         // We potentially have map information for this declaration already.
         // Look for the first set of components that refer to it.
         if (It != Info.end()) {
for (const auto *C :
         if (isa<MemberExpr>(IE)) {
           // Insert the pointer into Info to be processed by
           // generateInfoForComponentList. Because it is a member pointer
           // without a pointee, no entry will be generated for it, therefore
           // we need to generate one after the whole struct has been processed.
           // Nonetheless, generateInfoForComponentList must be called to take
           // the pointer into account for the calculation of the range of the
           // partial struct.
-          InfoGen(nullptr, L.second, OMPC_MAP_unknown, llvm::None,
-                  /*ReturnDevicePointer=*/false, C->isImplicit());
+          InfoGen(nullptr, Components, OMPC_MAP_unknown, llvm::None,
+                  /*ReturnDevicePointer=*/false, C->isImplicit(), nullptr);
           DeferredInfo[nullptr].emplace_back(IE, VD);
         } else {
           llvm::Value *Ptr =
               CGF.EmitLoadOfScalar(CGF.EmitLValue(IE), IE->getExprLoc());
           BasePointers.emplace_back(Ptr, VD);
           Pointers.push_back(Ptr);
           Sizes.push_back(llvm::Constant::getNullValue(CGF.Int64Ty));
           Types.push_back(OMP_MAP_RETURN_PARAM | OMP_MAP_TARGET_PARAM);
+          Mappers.push_back(nullptr);
         }
       }
     }
 
     for (const auto &M : Info) {
       // We need to know when we generate information for the first component
       // associated with a capture, because the mapping flags depend on it.
       bool IsFirstComponentList = true;
 
       // Temporary versions of arrays
       MapBaseValuesArrayTy CurBasePointers;
       MapValuesArrayTy CurPointers;
       MapValuesArrayTy CurSizes;
       MapFlagsArrayTy CurTypes;
+      MapMappersArrayTy CurMappers;
       StructRangeInfoTy PartialStruct;
 
       for (const MapInfo &L : M.second) {
         assert(!L.Components.empty() &&
                "Not expecting declaration with no component lists.");
 
         // Remember the current base pointer index.
         unsigned CurrentBasePointersIdx = CurBasePointers.size();
-        generateInfoForComponentList(L.MapType, L.MapModifiers, L.Components,
-                                     CurBasePointers, CurPointers, CurSizes,
-                                     CurTypes, PartialStruct,
-                                     IsFirstComponentList, L.IsImplicit);
+        generateInfoForComponentList(
+            L.MapType, L.MapModifiers, L.Components, CurBasePointers,
+            CurPointers, CurSizes, CurTypes, CurMappers, PartialStruct,
+            IsFirstComponentList, L.IsImplicit, L.Mapper);
 
         // If this entry relates with a device pointer, set the relevant
         // declaration and add the 'return pointer' flag.
         if (L.ReturnDevicePointer) {
           assert(CurBasePointers.size() > CurrentBasePointersIdx &&
                  "Unexpected number of mapped base pointers.");
 
           const ValueDecl *RelevantVD =
for (const auto &M : Info) {
           CurBasePointers.emplace_back(BasePtr, L.VD);
           CurPointers.push_back(Ptr);
           CurSizes.push_back(llvm::Constant::getNullValue(this->CGF.Int64Ty));
           // Entry is PTR_AND_OBJ and RETURN_PARAM. Also, set the placeholder
           // value MEMBER_OF=FFFF so that the entry is later updated with the
           // correct value of MEMBER_OF.
           CurTypes.push_back(OMP_MAP_PTR_AND_OBJ | OMP_MAP_RETURN_PARAM |
                              OMP_MAP_MEMBER_OF);
+          CurMappers.push_back(nullptr);
         }
       }
 
       // If there is an entry in PartialStruct it means we have a struct with
       // individual members mapped. Emit an extra combined entry.
       if (PartialStruct.Base.isValid())
-        emitCombinedEntry(BasePointers, Pointers, Sizes, Types, CurTypes,
-                          PartialStruct);
+        emitCombinedEntry(BasePointers, Pointers, Sizes, Types, Mappers,
+                          CurTypes, PartialStruct);
 
       // We need to append the results of this capture to what we already have.
       BasePointers.append(CurBasePointers.begin(), CurBasePointers.end());
       Pointers.append(CurPointers.begin(), CurPointers.end());
       Sizes.append(CurSizes.begin(), CurSizes.end());
       Types.append(CurTypes.begin(), CurTypes.end());
+      Mappers.append(CurMappers.begin(), CurMappers.end());
     }
   }
 
-  /// Generate all the base pointers, section pointers, sizes and map types for
-  /// the extracted map clauses of user-defined mapper.
+  /// Generate all the base pointers, section pointers, sizes, map types, and
+  /// mappers for the extracted map clauses of user-defined mapper.
   void generateAllInfoForMapper(MapBaseValuesArrayTy &BasePointers,
                                 MapValuesArrayTy &Pointers,
-                                MapValuesArrayTy &Sizes,
-                                MapFlagsArrayTy &Types) const {
+                                MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
+                                MapMappersArrayTy &Mappers) const {
     assert(CurDir.is<const OMPDeclareMapperDecl *>() &&
            "Expect a declare mapper directive");
     const auto *CurMapperDir = CurDir.get<const OMPDeclareMapperDecl *>();
     // We have to process the component lists that relate with the same
     // declaration in a single chunk so that we can generate the map flags
     // correctly. Therefore, we organize all lists in a map.
     llvm::MapVector<const ValueDecl *, SmallVector<MapInfo, 8>> Info;
 
-    // Helper function to fill the information map for the different supported
-    // clauses.
-    auto &&InfoGen = [&Info](
-        const ValueDecl *D,
-        OMPClauseMappableExprCommon::MappableExprComponentListRef L,
-        OpenMPMapClauseKind MapType,
-        ArrayRef<OpenMPMapModifierKind> MapModifiers,
-        bool ReturnDevicePointer, bool IsImplicit) {
-      const ValueDecl *VD =
-          D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
-      Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer,
-                            IsImplicit);
-    };
-
+    // Fill the information map for map clauses.
     for (const auto *C : CurMapperDir->clauselists()) {
-      const auto *MC = cast<OMPMapClause>(C);
+      const auto *MC = cast<const OMPMapClause>(C);
       for (const auto &L : MC->component_lists()) {
-        InfoGen(L.first, L.second, MC->getMapType(), MC->getMapTypeModifiers(),
-                /*ReturnDevicePointer=*/false, MC->isImplicit());
+        const ValueDecl *VD =
+            std::get<0>(L) ? cast<ValueDecl>(std::get<0>(L)->getCanonicalDecl())
+                           : nullptr;
+        // Get the corresponding user-defined mapper.
+        Info[VD].emplace_back(
+            std::get<1>(L), MC->getMapType(), MC->getMapTypeModifiers(),
+            /*ReturnDevicePointer=*/false, MC->isImplicit(), std::get<2>(L));
       }
     }
 
     for (const auto &M : Info) {
       // We need to know when we generate information for the first component
       // associated with a capture, because the mapping flags depend on it.
       bool IsFirstComponentList = true;
 
       // Temporary versions of arrays
       MapBaseValuesArrayTy CurBasePointers;
       MapValuesArrayTy CurPointers;
       MapValuesArrayTy CurSizes;
       MapFlagsArrayTy CurTypes;
+      MapMappersArrayTy CurMappers;
       StructRangeInfoTy PartialStruct;
 
       for (const MapInfo &L : M.second) {
         assert(!L.Components.empty() &&
                "Not expecting declaration with no component lists.");
-        generateInfoForComponentList(L.MapType, L.MapModifiers, L.Components,
-                                     CurBasePointers, CurPointers, CurSizes,
-                                     CurTypes, PartialStruct,
-                                     IsFirstComponentList, L.IsImplicit);
+        generateInfoForComponentList(
+            L.MapType, L.MapModifiers, L.Components, CurBasePointers,
+            CurPointers, CurSizes, CurTypes, CurMappers, PartialStruct,
+            IsFirstComponentList, L.IsImplicit, L.Mapper);
         IsFirstComponentList = false;
       }
 
       // If there is an entry in PartialStruct it means we have a struct with
       // individual members mapped. Emit an extra combined entry.
       if (PartialStruct.Base.isValid())
-        emitCombinedEntry(BasePointers, Pointers, Sizes, Types, CurTypes,
-                          PartialStruct);
+        emitCombinedEntry(BasePointers, Pointers, Sizes, Types, Mappers,
+                          CurTypes, PartialStruct);
 
       // We need to append the results of this capture to what we already have.
       BasePointers.append(CurBasePointers.begin(), CurBasePointers.end());
       Pointers.append(CurPointers.begin(), CurPointers.end());
       Sizes.append(CurSizes.begin(), CurSizes.end());
       Types.append(CurTypes.begin(), CurTypes.end());
+      Mappers.append(CurMappers.begin(), CurMappers.end());
     }
   }
 
   /// Emit capture info for lambdas for variables captured by reference.
   void generateInfoForLambdaCaptures(
       const ValueDecl *VD, llvm::Value *Arg, MapBaseValuesArrayTy &BasePointers,
       MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes,
-      MapFlagsArrayTy &Types,
+      MapFlagsArrayTy &Types, MapMappersArrayTy &Mappers,
       llvm::DenseMap<llvm::Value *, llvm::Value *> &LambdaPointers) const {
     const auto *RD = VD->getType()
                          .getCanonicalType()
                          .getNonReferenceType()
                          ->getAsCXXRecordDecl();
     if (!RD || !RD->isLambda())
       return;
     Address VDAddr = Address(Arg, CGF.getContext().getDeclAlign(VD));
if (ThisCapture) {
       LambdaPointers.try_emplace(ThisLVal.getPointer(), VDLVal.getPointer());
       BasePointers.push_back(ThisLVal.getPointer());
       Pointers.push_back(ThisLValVal.getPointer());
       Sizes.push_back(
           CGF.Builder.CreateIntCast(CGF.getTypeSize(CGF.getContext().VoidPtrTy),
                                     CGF.Int64Ty, /*isSigned=*/true));
       Types.push_back(OMP_MAP_PTR_AND_OBJ | OMP_MAP_LITERAL |
                       OMP_MAP_MEMBER_OF | OMP_MAP_IMPLICIT);
+      Mappers.push_back(nullptr);
     }
     for (const LambdaCapture &LC : RD->captures()) {
       if (!LC.capturesVariable())
         continue;
       const VarDecl *VD = LC.getCapturedVar();
       if (LC.getCaptureKind() != LCK_ByRef && !VD->getType()->isPointerType())
         continue;
       auto It = Captures.find(VD);
for (const LambdaCapture &LC : RD->captures()) {
         RValue VarRVal = CGF.EmitLoadOfLValue(VarLVal, RD->getLocation());
         LambdaPointers.try_emplace(VarLVal.getPointer(), VDLVal.getPointer());
         BasePointers.push_back(VarLVal.getPointer());
         Pointers.push_back(VarRVal.getScalarVal());
         Sizes.push_back(llvm::ConstantInt::get(CGF.Int64Ty, 0));
       }
       Types.push_back(OMP_MAP_PTR_AND_OBJ | OMP_MAP_LITERAL |
                       OMP_MAP_MEMBER_OF | OMP_MAP_IMPLICIT);
+      Mappers.push_back(nullptr);
     }
   }
 
   /// Set correct indices for lambdas captures.
   void adjustMemberOfForLambdaCaptures(
       const llvm::DenseMap<llvm::Value *, llvm::Value *> &LambdaPointers,
       MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers,
       MapFlagsArrayTy &Types) const {
for (unsigned I = 0, E = Types.size(); I < E; ++I) {
       // All other current entries will be MEMBER_OF the combined entry
       // (except for PTR_AND_OBJ entries which do not have a placeholder value
       // 0xFFFF in the MEMBER_OF field).
       OpenMPOffloadMappingFlags MemberOfFlag = getMemberOfFlag(TgtIdx);
       setCorrectMemberOfFlag(Types[I], MemberOfFlag);
     }
   }
 
-  /// Generate the base pointers, section pointers, sizes and map types
-  /// associated to a given capture.
+  /// Generate the base pointers, section pointers, sizes, map types, and
+  /// mappers associated to a given capture.
   void generateInfoForCapture(const CapturedStmt::Capture *Cap,
                               llvm::Value *Arg,
                               MapBaseValuesArrayTy &BasePointers,
                               MapValuesArrayTy &Pointers,
                               MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
+                              MapMappersArrayTy &Mappers,
                               StructRangeInfoTy &PartialStruct) const {
     assert(!Cap->capturesVariableArrayType() &&
            "Not expecting to generate map info for a variable array type!");
 
     // We need to know when we generating information for the first component
     const ValueDecl *VD = Cap->capturesThis()
                               ? nullptr
                               : Cap->getCapturedVar()->getCanonicalDecl();
 
     // If this declaration appears in a is_device_ptr clause we just have to
     // pass the pointer by value. If it is a reference to a declaration, we just
     // pass its value.
     if (DevPointersMap.count(VD)) {
       BasePointers.emplace_back(Arg, VD);
       Pointers.push_back(Arg);
       Sizes.push_back(
           CGF.Builder.CreateIntCast(CGF.getTypeSize(CGF.getContext().VoidPtrTy),
                                     CGF.Int64Ty, /*isSigned=*/true));
       Types.push_back(OMP_MAP_LITERAL | OMP_MAP_TARGET_PARAM);
+      Mappers.push_back(nullptr);
       return;
     }
 
     using MapData =
         std::tuple<OMPClauseMappableExprCommon::MappableExprComponentListRef,
-                   OpenMPMapClauseKind, ArrayRef<OpenMPMapModifierKind>, bool>;
+                   OpenMPMapClauseKind, ArrayRef<OpenMPMapModifierKind>, bool,
+                   const ValueDecl *>;
     SmallVector<MapData, 4> DeclComponentLists;
     assert(CurDir.is<const OMPExecutableDirective *>() &&
            "Expect a executable directive");
     const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
+    std::cerr << "C " << "\n";
     for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>()) {
       for (const auto &L : C->decl_component_lists(VD)) {
-        assert(L.first == VD &&
-               "We got information for the wrong declaration??");
-        assert(!L.second.empty() &&
+        const ValueDecl *VDecl, *Mapper;
+        OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
+        std::tie(VDecl, Components, Mapper) = L;
+        assert(VDecl == VD && "We got information for the wrong declaration??");
+        assert(!Components.empty() &&
                "Not expecting declaration with no component lists.");
-        DeclComponentLists.emplace_back(L.second, C->getMapType(),
+        DeclComponentLists.emplace_back(Components, C->getMapType(),
                                         C->getMapTypeModifiers(),
-                                        C->isImplicit());
+                                        C->isImplicit(), Mapper);
       }
     }
 
     // Find overlapping elements (including the offset from the base element).
     llvm::SmallDenseMap<
         const MapData *,
         llvm::SmallVector<
             OMPClauseMappableExprCommon::MappableExprComponentListRef, 4>,
         4>
         OverlappedData;
     size_t Count = 0;
     for (const MapData &L : DeclComponentLists) {
       OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
       OpenMPMapClauseKind MapType;
       ArrayRef<OpenMPMapModifierKind> MapModifiers;
       bool IsImplicit;
-      std::tie(Components, MapType, MapModifiers, IsImplicit) = L;
+      const ValueDecl *Mapper;
+      std::tie(Components, MapType, MapModifiers, IsImplicit, Mapper) = L;
       ++Count;
       for (const MapData &L1 : makeArrayRef(DeclComponentLists).slice(Count)) {
         OMPClauseMappableExprCommon::MappableExprComponentListRef Components1;
-        std::tie(Components1, MapType, MapModifiers, IsImplicit) = L1;
+        std::tie(Components1, MapType, MapModifiers, IsImplicit, Mapper) = L1;
         auto CI = Components.rbegin();
         auto CE = Components.rend();
         auto SI = Components1.rbegin();
         auto SE = Components1.rend();
         for (; CI != CE && SI != SE; ++CI, ++SI) {
           if (CI->getAssociatedExpression()->getStmtClass() !=
               SI->getAssociatedExpression()->getStmtClass())
             break;
void generateInfoForCapture(const CapturedStmt::Capture *Cap,
     // Associated with a capture, because the mapping flags depend on it.
     // Go through all of the elements with the overlapped elements.
     for (const auto &Pair : OverlappedData) {
       const MapData &L = *Pair.getFirst();
       OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
       OpenMPMapClauseKind MapType;
       ArrayRef<OpenMPMapModifierKind> MapModifiers;
       bool IsImplicit;
-      std::tie(Components, MapType, MapModifiers, IsImplicit) = L;
+      const ValueDecl *Mapper;
+      std::tie(Components, MapType, MapModifiers, IsImplicit, Mapper) = L;
       ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
           OverlappedComponents = Pair.getSecond();
       bool IsFirstComponentList = true;
       generateInfoForComponentList(MapType, MapModifiers, Components,
                                    BasePointers, Pointers, Sizes, Types,
-                                   PartialStruct, IsFirstComponentList,
-                                   IsImplicit, OverlappedComponents);
+                                   Mappers, PartialStruct, IsFirstComponentList,
+                                   IsImplicit, Mapper, OverlappedComponents);
     }
     // Go through other elements without overlapped elements.
     bool IsFirstComponentList = OverlappedData.empty();
     for (const MapData &L : DeclComponentLists) {
       OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
       OpenMPMapClauseKind MapType;
       ArrayRef<OpenMPMapModifierKind> MapModifiers;
       bool IsImplicit;
-      std::tie(Components, MapType, MapModifiers, IsImplicit) = L;
+      const ValueDecl *Mapper;
+      std::tie(Components, MapType, MapModifiers, IsImplicit, Mapper) = L;
       auto It = OverlappedData.find(&L);
       if (It == OverlappedData.end())
         generateInfoForComponentList(MapType, MapModifiers, Components,
                                      BasePointers, Pointers, Sizes, Types,
-                                     PartialStruct, IsFirstComponentList,
-                                     IsImplicit);
+                                     Mappers, PartialStruct,
+                                     IsFirstComponentList, IsImplicit, Mapper);
       IsFirstComponentList = false;
     }
   }
 
-  /// Generate the base pointers, section pointers, sizes and map types
-  /// associated with the declare target link variables.
+  /// Generate the base pointers, section pointers, sizes, map types, and
+  /// mappers associated with the declare target link variables.
   void generateInfoForDeclareTargetLink(MapBaseValuesArrayTy &BasePointers,
                                         MapValuesArrayTy &Pointers,
                                         MapValuesArrayTy &Sizes,
-                                        MapFlagsArrayTy &Types) const {
+                                        MapFlagsArrayTy &Types,
+                                        MapMappersArrayTy &Mappers) const {
     assert(CurDir.is<const OMPExecutableDirective *>() &&
            "Expect a executable directive");
     const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
     // Map other list items in the map clause which are not captured variables
     // but "declare target link" global variables.
     for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>()) {
       for (const auto &L : C->component_lists()) {
-        if (!L.first)
+        if (!std::get<0>(L))
           continue;
-        const auto *VD = dyn_cast<VarDecl>(L.first);
+        const auto *VD = dyn_cast<VarDecl>(std::get<0>(L));
         if (!VD)
           continue;
         llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
             OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
         if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
             !Res || *Res != OMPDeclareTargetDeclAttr::MT_Link)
           continue;
         StructRangeInfoTy PartialStruct;
         generateInfoForComponentList(
-            C->getMapType(), C->getMapTypeModifiers(), L.second, BasePointers,
-            Pointers, Sizes, Types, PartialStruct,
+            C->getMapType(), C->getMapTypeModifiers(), std::get<1>(L),
+            BasePointers, Pointers, Sizes, Types, Mappers, PartialStruct,
             /*IsFirstComponentList=*/true, C->isImplicit());
         assert(!PartialStruct.Base.isValid() &&
                "No partial structs for declare target link expected.");
       }
     }
   }
 
   /// Generate the default map information for a given capture \a CI,
   /// record field declaration \a RI and captured value \a CV.
   void generateDefaultMapInfo(const CapturedStmt::Capture &CI,
                               const FieldDecl &RI, llvm::Value *CV,
                               MapBaseValuesArrayTy &CurBasePointers,
                               MapValuesArrayTy &CurPointers,
                               MapValuesArrayTy &CurSizes,
-                              MapFlagsArrayTy &CurMapTypes) const {
+                              MapFlagsArrayTy &CurMapTypes,
+                              MapMappersArrayTy &CurMappers) const {
     bool IsImplicit = true;
     // Do the default mapping.
     if (CI.capturesThis()) {
       CurBasePointers.push_back(CV);
       CurPointers.push_back(CV);
       const auto *PtrTy = cast<PointerType>(RI.getType().getTypePtr());
       CurSizes.push_back(
           CGF.Builder.CreateIntCast(CGF.getTypeSize(PtrTy->getPointeeType()),
if (CI.capturesThis()) {
         IsImplicit = I->getSecond();
     }
     // Every default map produces a single argument which is a target parameter.
     CurMapTypes.back() |= OMP_MAP_TARGET_PARAM;
 
     // Add flag stating this is an implicit map.
     if (IsImplicit)
       CurMapTypes.back() |= OMP_MAP_IMPLICIT;
+
+    // No user-defined mapper for default mapping.
+    CurMappers.push_back(nullptr);
   }
 };
 } // anonymous namespace
 
 /// Emit the arrays used to pass the captures and map information to the
 /// offloading runtime library. If there is no map or capture information,
 /// return nullptr by reference.
 static void
 emitOffloadingArrays(CodeGenFunction &CGF,
                      MappableExprsHandler::MapBaseValuesArrayTy &BasePointers,
                      MappableExprsHandler::MapValuesArrayTy &Pointers,
                      MappableExprsHandler::MapValuesArrayTy &Sizes,
                      MappableExprsHandler::MapFlagsArrayTy &MapTypes,
+                     MappableExprsHandler::MapMappersArrayTy &Mappers,
                      CGOpenMPRuntime::TargetDataInfo &Info) {
   CodeGenModule &CGM = CGF.CGM;
   ASTContext &Ctx = CGF.getContext();
 
   // Reset the array information.
   Info.clearArrayInfo();
   Info.NumberOfPtrs = BasePointers.size();
 
if (Info.NumberOfPtrs) {
     QualType PointerArrayType =
         Ctx.getConstantArrayType(Ctx.VoidPtrTy, PointerNumAP, ArrayType::Normal,
                                  /*IndexTypeQuals=*/0);
 
     Info.BasePointersArray =
         CGF.CreateMemTemp(PointerArrayType, ".offload_baseptrs").getPointer();
     Info.PointersArray =
         CGF.CreateMemTemp(PointerArrayType, ".offload_ptrs").getPointer();
+    Info.MappersArray =
+        CGF.CreateMemTemp(PointerArrayType, ".offload_mappers").getPointer();
 
     // If we don't have any VLA types or other types that require runtime
     // evaluation, we can use a constant array for the map sizes, otherwise we
     // need to fill up the arrays as we do for the pointers.
     QualType Int64Ty =
         Ctx.getIntTypeForBitwidth(/*DestWidth=*/64, /*Signed=*/1);
     if (hasRuntimeEvaluationCaptureSize) {
       QualType SizeArrayType =
for (unsigned I = 0; I < Info.NumberOfPtrs; ++I) {
             Info.SizesArray,
             /*Idx0=*/0,
             /*Idx1=*/I);
         Address SAddr(S, Ctx.getTypeAlignInChars(Int64Ty));
         CGF.Builder.CreateStore(
             CGF.Builder.CreateIntCast(Sizes[I], CGM.Int64Ty, /*isSigned=*/true),
             SAddr);
       }
+
+      // Fill up the mapper array.
+      llvm::Value *MFunc = llvm::ConstantPointerNull::get(CGM.VoidPtrTy);
+      if (Mappers[I])
+        MFunc = CGM.getOpenMPRuntime().getUserDefinedMapperFunc(
+            cast<OMPDeclareMapperDecl>(Mappers[I]));
+      llvm::Value *M = CGF.Builder.CreateConstInBoundsGEP2_32(
+          llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
+          Info.MappersArray, 0, I);
+      M = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
+          M, MFunc->getType()->getPointerTo(/*AddrSpace=*/0));
+      Address MAddr(M, Ctx.getTypeAlignInChars(Ctx.VoidPtrTy));
+      CGF.Builder.CreateStore(MFunc, MAddr);
     }
   }
 }
 
 /// Emit the arguments to be passed to the runtime library based on the
-/// arrays of pointers, sizes and map types.
+/// arrays of base pointers, pointers, sizes, map types, and mappers.
 static void emitOffloadingArraysArgument(
     CodeGenFunction &CGF, llvm::Value *&BasePointersArrayArg,
     llvm::Value *&PointersArrayArg, llvm::Value *&SizesArrayArg,
-    llvm::Value *&MapTypesArrayArg, CGOpenMPRuntime::TargetDataInfo &Info) {
+    llvm::Value *&MapTypesArrayArg, llvm::Value *&MappersArrayArg,
+    CGOpenMPRuntime::TargetDataInfo &Info) {
   CodeGenModule &CGM = CGF.CGM;
   if (Info.NumberOfPtrs) {
     BasePointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
         llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
         Info.BasePointersArray,
         /*Idx0=*/0, /*Idx1=*/0);
     PointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
         llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
         Info.PointersArray,
         /*Idx0=*/0,
         /*Idx1=*/0);
     SizesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
         llvm::ArrayType::get(CGM.Int64Ty, Info.NumberOfPtrs), Info.SizesArray,
         /*Idx0=*/0, /*Idx1=*/0);
     MapTypesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
         llvm::ArrayType::get(CGM.Int64Ty, Info.NumberOfPtrs),
         Info.MapTypesArray,
         /*Idx0=*/0,
         /*Idx1=*/0);
+    MappersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
+        llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
+        Info.MappersArray,
+        /*Idx0=*/0, /*Idx1=*/0);
   } else {
     BasePointersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
     PointersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
     SizesArrayArg = llvm::ConstantPointerNull::get(CGM.Int64Ty->getPointerTo());
     MapTypesArrayArg =
         llvm::ConstantPointerNull::get(CGM.Int64Ty->getPointerTo());
+    MappersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
   }
 }
 
 /// Check for inner distribute directive.
 static const OMPExecutableDirective *
 getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) {
   const auto *CS = D.getInnermostCapturedStmt();
   const auto *Body =
void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D,
   // Evaluate whether the initial condition is satisfied.
   llvm::Value *IsEmpty =
       MapperCGF.Builder.CreateICmpEQ(PtrBegin, PtrEnd, "omp.arraymap.isempty");
   MapperCGF.Builder.CreateCondBr(IsEmpty, DoneBB, BodyBB);
   llvm::BasicBlock *EntryBB = MapperCGF.Builder.GetInsertBlock();
 
   // Emit the loop body block.
   MapperCGF.EmitBlock(BodyBB);
+  llvm::BasicBlock *LastBB = BodyBB;
   llvm::PHINode *PtrPHI = MapperCGF.Builder.CreatePHI(
       PtrBegin->getType(), 2, "omp.arraymap.ptrcurrent");
   PtrPHI->addIncoming(PtrBegin, EntryBB);
   Address PtrCurrent =
       Address(PtrPHI, MapperCGF.GetAddrOfLocalVar(&BeginArg)
                           .getAlignment()
                           .alignmentOfArrayElement(ElementSize));
   // Privatize the declared variable of mapper to be the current array element.
   CodeGenFunction::OMPPrivateScope Scope(MapperCGF);
   Scope.addPrivate(MapperVarDecl, [&MapperCGF, PtrCurrent, PtrTy]() {
     return MapperCGF
         .EmitLoadOfPointerLValue(PtrCurrent, PtrTy->castAs<PointerType>())
         .getAddress();
   });
   (void)Scope.Privatize();
 
   // Get map clause information. Fill up the arrays with all mapped variables.
   MappableExprsHandler::MapBaseValuesArrayTy BasePointers;
   MappableExprsHandler::MapValuesArrayTy Pointers;
   MappableExprsHandler::MapValuesArrayTy Sizes;
   MappableExprsHandler::MapFlagsArrayTy MapTypes;
+  MappableExprsHandler::MapMappersArrayTy Mappers;
   MappableExprsHandler MEHandler(*D, MapperCGF);
-  MEHandler.generateAllInfoForMapper(BasePointers, Pointers, Sizes, MapTypes);
+  MEHandler.generateAllInfoForMapper(BasePointers, Pointers, Sizes, MapTypes,
+                                     Mappers);
 
   // Call the runtime API __tgt_mapper_num_components to get the number of
   // pre-existing components.
   llvm::Value *OffloadingArgs[] = {Handle};
   llvm::Value *PreviousSize = MapperCGF.EmitRuntimeCall(
       createRuntimeFunction(OMPRTL__tgt_mapper_num_components), OffloadingArgs);
   llvm::Value *ShiftedPreviousSize = MapperCGF.Builder.CreateShl(
       PreviousSize,
for (unsigned I = 0; I < BasePointers.size(); ++I) {
     MapperCGF.Builder.CreateCondBr(IsFrom, FromBB, EndBB);
     // In case of from, clear OMP_MAP_TO.
     MapperCGF.EmitBlock(FromBB);
     llvm::Value *FromMapType = MapperCGF.Builder.CreateAnd(
         MemberMapType,
         MapperCGF.Builder.getInt64(~MappableExprsHandler::OMP_MAP_TO));
     // In case of tofrom, do nothing.
     MapperCGF.EmitBlock(EndBB);
+    LastBB = EndBB;
     llvm::PHINode *CurMapType =
         MapperCGF.Builder.CreatePHI(CGM.Int64Ty, 4, "omp.maptype");
     CurMapType->addIncoming(AllocMapType, AllocBB);
     CurMapType->addIncoming(ToMapType, ToBB);
     CurMapType->addIncoming(FromMapType, FromBB);
     CurMapType->addIncoming(MemberMapType, ToElseBB);
 
-    // TODO: call the corresponding mapper function if a user-defined mapper is
-    // associated with this map clause.
-    // Call the runtime API __tgt_push_mapper_component to fill up the runtime
-    // data structure.
     llvm::Value *OffloadingArgs[] = {Handle, CurBaseArg, CurBeginArg,
                                      CurSizeArg, CurMapType};
+    if (Mappers[I]) {
+      // Call the corresponding mapper function.
+      llvm::Function *MapperFunc =
+          getUserDefinedMapperFunc(cast<OMPDeclareMapperDecl>(Mappers[I]));
+      assert(MapperFunc && "Expect a valid mapper function is available.");
+      MapperCGF.Builder.CreateCall(MapperFunc, OffloadingArgs);
+    } else {
+      // Call the runtime API __tgt_push_mapper_component to fill up the runtime
+      // data structure.
       MapperCGF.EmitRuntimeCall(
           createRuntimeFunction(OMPRTL__tgt_push_mapper_component),
           OffloadingArgs);
     }
+  }
 
   // Update the pointer to point to the next element that needs to be mapped,
   // and check whether we have mapped all elements.
   llvm::Value *PtrNext = MapperCGF.Builder.CreateConstGEP1_32(
       PtrPHI, /*Idx0=*/1, "omp.arraymap.next");
-  PtrPHI->addIncoming(PtrNext, BodyBB);
+  PtrPHI->addIncoming(PtrNext, LastBB);
   llvm::Value *IsDone =
       MapperCGF.Builder.CreateICmpEQ(PtrNext, PtrEnd, "omp.arraymap.isdone");
   llvm::BasicBlock *ExitBB = MapperCGF.createBasicBlock("omp.arraymap.exit");
   MapperCGF.Builder.CreateCondBr(IsDone, ExitBB, BodyBB);
 
   MapperCGF.EmitBlock(ExitBB);
   // Emit array deletion if this is an array section and \p MapType indicates
   // that deletion is required.
llvm::Value *MapTypeArg = MapperCGF.Builder.CreateAnd(
                                    MappableExprsHandler::OMP_MAP_FROM)));
   // Call the runtime API __tgt_push_mapper_component to fill up the runtime
   // data structure.
   llvm::Value *OffloadingArgs[] = {Handle, Base, Begin, ArraySize, MapTypeArg};
   MapperCGF.EmitRuntimeCall(
       createRuntimeFunction(OMPRTL__tgt_push_mapper_component), OffloadingArgs);
 }
 
+llvm::Function *
+CGOpenMPRuntime::getUserDefinedMapperFunc(const OMPDeclareMapperDecl *D) {
+  auto I = UDMMap.find(D);
+  if (I != UDMMap.end())
+    return I->second;
+  emitUserDefinedMapper(D);
+  return UDMMap.lookup(D);
+}
+
 void CGOpenMPRuntime::emitTargetNumIterationsCall(
     CodeGenFunction &CGF, const OMPExecutableDirective &D, const Expr *Device,
     const llvm::function_ref<llvm::Value *(
         CodeGenFunction &CGF, const OMPLoopDirective &D)> &SizeEmitter) {
   OpenMPDirectiveKind Kind = D.getDirectiveKind();
   const OMPExecutableDirective *TD = &D;
   // Get nested teams distribute kind directive, if any.
   if (!isOpenMPDistributeDirective(Kind) || !isOpenMPTeamsDirective(Kind))
if (NumTeams) {
                            "with number of teams.");
       llvm::Value *OffloadingArgs[] = {DeviceID,
                                        OutlinedFnID,
                                        PointerNum,
                                        InputInfo.BasePointersArray.getPointer(),
                                        InputInfo.PointersArray.getPointer(),
                                        InputInfo.SizesArray.getPointer(),
                                        MapTypesArray,
+                                       InputInfo.MappersArray.getPointer(),
                                        NumTeams,
                                        NumThreads};
       Return = CGF.EmitRuntimeCall(
-          createRuntimeFunction(HasNowait ? OMPRTL__tgt_target_teams_nowait
-                                          : OMPRTL__tgt_target_teams),
+          createRuntimeFunction(HasNowait
+                                    ? OMPRTL__tgt_target_teams_nowait_mapper
+                                    : OMPRTL__tgt_target_teams_mapper),
           OffloadingArgs);
     } else {
       llvm::Value *OffloadingArgs[] = {DeviceID,
                                        OutlinedFnID,
                                        PointerNum,
                                        InputInfo.BasePointersArray.getPointer(),
                                        InputInfo.PointersArray.getPointer(),
                                        InputInfo.SizesArray.getPointer(),
-                                       MapTypesArray};
+                                       MapTypesArray,
+                                       InputInfo.MappersArray.getPointer()};
       Return = CGF.EmitRuntimeCall(
-          createRuntimeFunction(HasNowait ? OMPRTL__tgt_target_nowait
-                                          : OMPRTL__tgt_target),
+          createRuntimeFunction(HasNowait ? OMPRTL__tgt_target_nowait_mapper
+                                          : OMPRTL__tgt_target_mapper),
           OffloadingArgs);
     }
 
     // Check the error code and execute the host version if required.
     llvm::BasicBlock *OffloadFailedBlock =
         CGF.createBasicBlock("omp_offload.failed");
     llvm::BasicBlock *OffloadContBlock =
         CGF.createBasicBlock("omp_offload.cont");
void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
   auto &&TargetThenGen = [this, &ThenGen, &D, &InputInfo, &MapTypesArray,
                           &CapturedVars, RequiresOuterTask,
                           &CS](CodeGenFunction &CGF, PrePostActionTy &) {
     // Fill up the arrays with all the captured variables.
     MappableExprsHandler::MapBaseValuesArrayTy BasePointers;
     MappableExprsHandler::MapValuesArrayTy Pointers;
     MappableExprsHandler::MapValuesArrayTy Sizes;
     MappableExprsHandler::MapFlagsArrayTy MapTypes;
+    MappableExprsHandler::MapMappersArrayTy Mappers;
 
     // Get mappable expression information.
     MappableExprsHandler MEHandler(D, CGF);
     llvm::DenseMap<llvm::Value *, llvm::Value *> LambdaPointers;
 
     auto RI = CS.getCapturedRecordDecl()->field_begin();
     auto CV = CapturedVars.begin();
     for (CapturedStmt::const_capture_iterator CI = CS.capture_begin(),
                                               CE = CS.capture_end();
          CI != CE; ++CI, ++RI, ++CV) {
       MappableExprsHandler::MapBaseValuesArrayTy CurBasePointers;
       MappableExprsHandler::MapValuesArrayTy CurPointers;
       MappableExprsHandler::MapValuesArrayTy CurSizes;
       MappableExprsHandler::MapFlagsArrayTy CurMapTypes;
+      MappableExprsHandler::MapMappersArrayTy CurMappers;
       MappableExprsHandler::StructRangeInfoTy PartialStruct;
+      (*CV)->dump();
 
       // VLA sizes are passed to the outlined region by copy and do not have map
       // information associated.
       if (CI->capturesVariableArrayType()) {
         CurBasePointers.push_back(*CV);
         CurPointers.push_back(*CV);
         CurSizes.push_back(CGF.Builder.CreateIntCast(
             CGF.getTypeSize(RI->getType()), CGF.Int64Ty, /*isSigned=*/true));
         // Copy to the device as an argument. No need to retrieve it.
         CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_LITERAL |
                               MappableExprsHandler::OMP_MAP_TARGET_PARAM |
                               MappableExprsHandler::OMP_MAP_IMPLICIT);
+        CurMappers.push_back(nullptr);
       } else {
         // If we have any information in the map clause, we use it, otherwise we
         // just do a default mapping.
         MEHandler.generateInfoForCapture(CI, *CV, CurBasePointers, CurPointers,
-                                         CurSizes, CurMapTypes, PartialStruct);
+                                         CurSizes, CurMapTypes, CurMappers,
+                                         PartialStruct);
         if (CurBasePointers.empty())
           MEHandler.generateDefaultMapInfo(*CI, **RI, *CV, CurBasePointers,
-                                           CurPointers, CurSizes, CurMapTypes);
+                                           CurPointers, CurSizes, CurMapTypes,
+                                           CurMappers);
         // Generate correct mapping for variables captured by reference in
         // lambdas.
         if (CI->capturesVariable())
           MEHandler.generateInfoForLambdaCaptures(
               CI->getCapturedVar(), *CV, CurBasePointers, CurPointers, CurSizes,
-              CurMapTypes, LambdaPointers);
+              CurMapTypes, CurMappers, LambdaPointers);
       }
       // We expect to have at least an element of information for this capture.
       assert(!CurBasePointers.empty() &&
              "Non-existing map pointer for capture!");
       assert(CurBasePointers.size() == CurPointers.size() &&
              CurBasePointers.size() == CurSizes.size() &&
              CurBasePointers.size() == CurMapTypes.size() &&
+             CurBasePointers.size() == CurMappers.size() &&
              "Inconsistent map information sizes!");
 
       // If there is an entry in PartialStruct it means we have a struct with
       // individual members mapped. Emit an extra combined entry.
       if (PartialStruct.Base.isValid())
         MEHandler.emitCombinedEntry(BasePointers, Pointers, Sizes, MapTypes,
-                                    CurMapTypes, PartialStruct);
+                                    Mappers, CurMapTypes, PartialStruct);
 
       // We need to append the results of this capture to what we already have.
       BasePointers.append(CurBasePointers.begin(), CurBasePointers.end());
       Pointers.append(CurPointers.begin(), CurPointers.end());
       Sizes.append(CurSizes.begin(), CurSizes.end());
       MapTypes.append(CurMapTypes.begin(), CurMapTypes.end());
+      Mappers.append(CurMappers.begin(), CurMappers.end());
     }
     // Adjust MEMBER_OF flags for the lambdas captures.
     MEHandler.adjustMemberOfForLambdaCaptures(LambdaPointers, BasePointers,
                                               Pointers, MapTypes);
     // Map other list items in the map clause which are not captured variables
     // but "declare target link" global variables.
     MEHandler.generateInfoForDeclareTargetLink(BasePointers, Pointers, Sizes,
-                                               MapTypes);
+                                               MapTypes, Mappers);
 
     TargetDataInfo Info;
     // Fill up the arrays and create the arguments.
-    emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+    emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Mappers,
+                         Info);
     emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
                                  Info.PointersArray, Info.SizesArray,
-                                 Info.MapTypesArray, Info);
+                                 Info.MapTypesArray, Info.MappersArray, Info);
     InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;
     InputInfo.BasePointersArray =
         Address(Info.BasePointersArray, CGM.getPointerAlign());
     InputInfo.PointersArray =
         Address(Info.PointersArray, CGM.getPointerAlign());
     InputInfo.SizesArray = Address(Info.SizesArray, CGM.getPointerAlign());
+    InputInfo.MappersArray = Address(Info.MappersArray, CGM.getPointerAlign());
     MapTypesArray = Info.MapTypesArray;
     if (RequiresOuterTask)
       CGF.EmitOMPTargetTaskBasedDirective(D, ThenGen, InputInfo);
     else
       emitInlinedDirective(CGF, D.getDirectiveKind(), ThenGen);
   };
 
   auto &&TargetElseGen = [this, &ElseGen, &D, RequiresOuterTask](
void CGOpenMPRuntime::emitTargetDataCalls(
   // closing of the region.
   auto &&BeginThenGen = [this, &D, Device, &Info,
                          &CodeGen](CodeGenFunction &CGF, PrePostActionTy &) {
     // Fill up the arrays with all the mapped variables.
     MappableExprsHandler::MapBaseValuesArrayTy BasePointers;
     MappableExprsHandler::MapValuesArrayTy Pointers;
     MappableExprsHandler::MapValuesArrayTy Sizes;
     MappableExprsHandler::MapFlagsArrayTy MapTypes;
+    MappableExprsHandler::MapMappersArrayTy Mappers;
 
     // Get map clause information.
-    MappableExprsHandler MCHandler(D, CGF);
-    MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);
+    MappableExprsHandler MEHandler(D, CGF);
+    MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes, Mappers);
 
     // Fill up the arrays and create the arguments.
-    emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+    emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Mappers,
+                         Info);
 
     llvm::Value *BasePointersArrayArg = nullptr;
     llvm::Value *PointersArrayArg = nullptr;
     llvm::Value *SizesArrayArg = nullptr;
     llvm::Value *MapTypesArrayArg = nullptr;
+    llvm::Value *MappersArrayArg = nullptr;
     emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
-                                 SizesArrayArg, MapTypesArrayArg, Info);
+                                 SizesArrayArg, MapTypesArrayArg,
+                                 MappersArrayArg, Info);
 
     // Emit device ID if any.
     llvm::Value *DeviceID = nullptr;
     if (Device) {
       DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
                                            CGF.Int64Ty, /*isSigned=*/true);
     } else {
       DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);
     }
 
     // Emit the number of elements in the offloading arrays.
     llvm::Value *PointerNum = CGF.Builder.getInt32(Info.NumberOfPtrs);
 
     llvm::Value *OffloadingArgs[] = {
-        DeviceID,         PointerNum,    BasePointersArrayArg,
-        PointersArrayArg, SizesArrayArg, MapTypesArrayArg};
-    CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_target_data_begin),
-                        OffloadingArgs);
+        DeviceID,      PointerNum,       BasePointersArrayArg, PointersArrayArg,
+        SizesArrayArg, MapTypesArrayArg, MappersArrayArg};
+    CGF.EmitRuntimeCall(
+        createRuntimeFunction(OMPRTL__tgt_target_data_begin_mapper),
 
     // If device pointer privatization is required, emit the body of the region
     // here. It will have to be duplicated: with and without privatization.
     if (!Info.CaptureDeviceAddrMap.empty())
       CodeGen(CGF);
   };
 
   // Generate code for the closing of the data region.
   auto &&EndThenGen = [this, Device, &Info](CodeGenFunction &CGF,
                                             PrePostActionTy &) {
     assert(Info.isValid() && "Invalid data environment closing arguments.");
 
     llvm::Value *BasePointersArrayArg = nullptr;
     llvm::Value *PointersArrayArg = nullptr;
     llvm::Value *SizesArrayArg = nullptr;
     llvm::Value *MapTypesArrayArg = nullptr;
+    llvm::Value *MappersArrayArg = nullptr;
     emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
-                                 SizesArrayArg, MapTypesArrayArg, Info);
+                                 SizesArrayArg, MapTypesArrayArg,
+                                 MappersArrayArg, Info);
 
     // Emit device ID if any.
     llvm::Value *DeviceID = nullptr;
     if (Device) {
       DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
                                            CGF.Int64Ty, /*isSigned=*/true);
     } else {
       DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);
     }
 
     // Emit the number of elements in the offloading arrays.
     llvm::Value *PointerNum = CGF.Builder.getInt32(Info.NumberOfPtrs);
 
     llvm::Value *OffloadingArgs[] = {
-        DeviceID,         PointerNum,    BasePointersArrayArg,
-        PointersArrayArg, SizesArrayArg, MapTypesArrayArg};
-    CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_target_data_end),
-                        OffloadingArgs);
+        DeviceID,      PointerNum,       BasePointersArrayArg, PointersArrayArg,
+        SizesArrayArg, MapTypesArrayArg, MappersArrayArg};
+    CGF.EmitRuntimeCall(
+        createRuntimeFunction(OMPRTL__tgt_target_data_end_mapper),
   };
 
   // If we need device pointer privatization, we need to emit the body of the
   // region with no privatization in the 'else' branch of the conditional.
   // Otherwise, we don't have to do anything.
   auto &&BeginElseGen = [&Info, &CodeGen, &NoPrivAction](CodeGenFunction &CGF,
                                                          PrePostActionTy &) {
     if (!Info.CaptureDeviceAddrMap.empty()) {
auto &&ThenGen = [this, &D, Device, &InputInfo,
     llvm::Constant *PointerNum =
         CGF.Builder.getInt32(InputInfo.NumberOfTargetItems);
 
     llvm::Value *OffloadingArgs[] = {DeviceID,
                                      PointerNum,
                                      InputInfo.BasePointersArray.getPointer(),
                                      InputInfo.PointersArray.getPointer(),
                                      InputInfo.SizesArray.getPointer(),
-                                     MapTypesArray};
+                                     MapTypesArray,
+                                     InputInfo.MappersArray.getPointer()};
 
-    // Select the right runtime function call for each expected standalone
+    // Select the right runtime function call for each standalone
     // directive.
     const bool HasNowait = D.hasClausesOfKind<OMPNowaitClause>();
     OpenMPRTLFunction RTLFn;
     switch (D.getDirectiveKind()) {
     case OMPD_target_enter_data:
-      RTLFn = HasNowait ? OMPRTL__tgt_target_data_begin_nowait
-                        : OMPRTL__tgt_target_data_begin;
+      RTLFn = HasNowait ? OMPRTL__tgt_target_data_begin_nowait_mapper
+                        : OMPRTL__tgt_target_data_begin_mapper;
       break;
     case OMPD_target_exit_data:
-      RTLFn = HasNowait ? OMPRTL__tgt_target_data_end_nowait
-                        : OMPRTL__tgt_target_data_end;
+      RTLFn = HasNowait ? OMPRTL__tgt_target_data_end_nowait_mapper
+                        : OMPRTL__tgt_target_data_end_mapper;
       break;
     case OMPD_target_update:
-      RTLFn = HasNowait ? OMPRTL__tgt_target_data_update_nowait
-                        : OMPRTL__tgt_target_data_update;
+      RTLFn = HasNowait ? OMPRTL__tgt_target_data_update_nowait_mapper
+                        : OMPRTL__tgt_target_data_update_mapper;
       break;
     case OMPD_parallel:
     case OMPD_for:
     case OMPD_parallel_for:
     case OMPD_parallel_sections:
     case OMPD_for_simd:
     case OMPD_parallel_for_simd:
     case OMPD_cancel:
void CGOpenMPRuntime::emitTargetDataStandAloneCall(
 
   auto &&TargetThenGen = [this, &ThenGen, &D, &InputInfo, &MapTypesArray](
                              CodeGenFunction &CGF, PrePostActionTy &) {
     // Fill up the arrays with all the mapped variables.
     MappableExprsHandler::MapBaseValuesArrayTy BasePointers;
     MappableExprsHandler::MapValuesArrayTy Pointers;
     MappableExprsHandler::MapValuesArrayTy Sizes;
     MappableExprsHandler::MapFlagsArrayTy MapTypes;
+    MappableExprsHandler::MapMappersArrayTy Mappers;
 
     // Get map clause information.
     MappableExprsHandler MEHandler(D, CGF);
-    MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);
+    MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes, Mappers);
 
     TargetDataInfo Info;
     // Fill up the arrays and create the arguments.
-    emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info);
+    emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Mappers,
+                         Info);
     emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
                                  Info.PointersArray, Info.SizesArray,
-                                 Info.MapTypesArray, Info);
+                                 Info.MapTypesArray, Info.MappersArray, Info);
     InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;
     InputInfo.BasePointersArray =
         Address(Info.BasePointersArray, CGM.getPointerAlign());
     InputInfo.PointersArray =
         Address(Info.PointersArray, CGM.getPointerAlign());
     InputInfo.SizesArray =
         Address(Info.SizesArray, CGM.getPointerAlign());
+    InputInfo.MappersArray = Address(Info.MappersArray, CGM.getPointerAlign());
     MapTypesArray = Info.MapTypesArray;
     if (D.hasClausesOfKind<OMPDependClause>())
       CGF.EmitOMPTargetTaskBasedDirective(D, ThenGen, InputInfo);
     else
       emitInlinedDirective(CGF, D.getDirectiveKind(), ThenGen);
   };
 
   if (IfCond) {

lib/CodeGen/CGStmtOpenMP.cpp

for (auto *IInit : C->private_copies()) {
       ++IRef;
       ++IElemInitRef;
     }
   }
   OMPPrivateScope TargetScope(*this);
   VarDecl *BPVD = nullptr;
   VarDecl *PVD = nullptr;
   VarDecl *SVD = nullptr;
+  VarDecl *MVD = nullptr;
   if (InputInfo.NumberOfTargetItems > 0) {
     auto *CD = CapturedDecl::Create(
         getContext(), getContext().getTranslationUnitDecl(), /*NumParams=*/0);
     llvm::APInt ArrSize(/*numBits=*/32, InputInfo.NumberOfTargetItems);
-    QualType BaseAndPointersType = getContext().getConstantArrayType(
+    QualType BaseAndPointerAndMapperType = getContext().getConstantArrayType(
         getContext().VoidPtrTy, ArrSize, ArrayType::Normal,
         /*IndexTypeQuals=*/0);
     BPVD = createImplicitFirstprivateForType(
-        getContext(), Data, BaseAndPointersType, CD, S.getBeginLoc());
+        getContext(), Data, BaseAndPointerAndMapperType, CD, S.getBeginLoc());
     PVD = createImplicitFirstprivateForType(
-        getContext(), Data, BaseAndPointersType, CD, S.getBeginLoc());
+        getContext(), Data, BaseAndPointerAndMapperType, CD, S.getBeginLoc());
     QualType SizesType = getContext().getConstantArrayType(
         getContext().getIntTypeForBitwidth(/*DestWidth=*/64, /*Signed=*/1),
         ArrSize, ArrayType::Normal,
         /*IndexTypeQuals=*/0);
     SVD = createImplicitFirstprivateForType(getContext(), Data, SizesType, CD,
                                             S.getBeginLoc());
+    MVD = createImplicitFirstprivateForType(
+        getContext(), Data, BaseAndPointerAndMapperType, CD, S.getBeginLoc());
     TargetScope.addPrivate(
         BPVD, [&InputInfo]() { return InputInfo.BasePointersArray; });
     TargetScope.addPrivate(PVD,
                            [&InputInfo]() { return InputInfo.PointersArray; });
     TargetScope.addPrivate(SVD,
                            [&InputInfo]() { return InputInfo.SizesArray; });
+    TargetScope.addPrivate(MVD,
+                           [&InputInfo]() { return InputInfo.MappersArray; });
   }
   (void)TargetScope.Privatize();
   // Build list of dependences.
   for (const auto *C : S.getClausesOfKind<OMPDependClause>())
     for (const Expr *IRef : C->varlists())
       Data.Dependences.emplace_back(C->getDependencyKind(), IRef);
-  auto &&CodeGen = [&Data, &S, CS, &BodyGen, BPVD, PVD, SVD,
+  auto &&CodeGen = [&Data, &S, CS, &BodyGen, BPVD, PVD, SVD, MVD,
                     &InputInfo](CodeGenFunction &CGF, PrePostActionTy &Action) {
     // Set proper addresses for generated private copies.
     OMPPrivateScope Scope(CGF);
     if (!Data.FirstprivateVars.empty()) {
       llvm::FunctionType *CopyFnTy = llvm::FunctionType::get(
           CGF.Builder.getVoidTy(), {CGF.Builder.getInt8PtrTy()}, true);
       enum { PrivatesParam = 2, CopyFnParam = 3 };
       llvm::Value *CopyFn = CGF.Builder.CreateLoad(
auto &&CodeGen = [&Data, &S, CS, &BodyGen, BPVD, PVD, SVD, MVD,
     (void)Scope.Privatize();
     if (InputInfo.NumberOfTargetItems > 0) {
       InputInfo.BasePointersArray = CGF.Builder.CreateConstArrayGEP(
           CGF.GetAddrOfLocalVar(BPVD), /*Index=*/0);
       InputInfo.PointersArray = CGF.Builder.CreateConstArrayGEP(
           CGF.GetAddrOfLocalVar(PVD), /*Index=*/0);
       InputInfo.SizesArray = CGF.Builder.CreateConstArrayGEP(
           CGF.GetAddrOfLocalVar(SVD), /*Index=*/0);
+      InputInfo.MappersArray = CGF.Builder.CreateConstArrayGEP(
+          CGF.GetAddrOfLocalVar(MVD), /*Index=*/0);
     }
 
     Action.Enter(CGF);
     OMPLexicalScope LexScope(CGF, S, OMPD_task, /*EmitPreInitStmt=*/false);
     BodyGen(CGF);
   };
   llvm::Function *OutlinedFn = CGM.getOpenMPRuntime().emitTaskOutlinedFunction(
       S, *I, *PartId, *TaskT, S.getDirectiveKind(), CodeGen, /*Tied=*/true,

lib/CodeGen/CodeGenFunction.h

   void EmitOMPTaskBasedDirective(const OMPExecutableDirective &S,
                                  const OpenMPDirectiveKind CapturedRegion,
                                  const RegionCodeGenTy &BodyGen,
                                  const TaskGenTy &TaskGen, OMPTaskDataTy &Data);
   struct OMPTargetDataInfo {
     Address BasePointersArray = Address::invalid();
     Address PointersArray = Address::invalid();
     Address SizesArray = Address::invalid();
+    Address MappersArray = Address::invalid();
     unsigned NumberOfTargetItems = 0;
     explicit OMPTargetDataInfo() = default;
     OMPTargetDataInfo(Address BasePointersArray, Address PointersArray,
-                      Address SizesArray, unsigned NumberOfTargetItems)
+                      Address SizesArray, Address MappersArray,
+                      unsigned NumberOfTargetItems)
         : BasePointersArray(BasePointersArray), PointersArray(PointersArray),
-          SizesArray(SizesArray), NumberOfTargetItems(NumberOfTargetItems) {}
+          SizesArray(SizesArray), MappersArray(MappersArray),
+          NumberOfTargetItems(NumberOfTargetItems) {}
   };
   void EmitOMPTargetTaskBasedDirective(const OMPExecutableDirective &S,
                                        const RegionCodeGenTy &BodyGen,
                                        OMPTargetDataInfo &InputInfo);
 
   void EmitOMPParallelDirective(const OMPParallelDirective &S);
   void EmitOMPSimdDirective(const OMPSimdDirective &S);
   void EmitOMPForDirective(const OMPForDirective &S);

test/OpenMP/capturing_in_templates.cpp

 
 template <typename T1, typename T2>
 pair<T1, T2> make_pair(T1 &&t1, T2 &&t2) {
   return {t1, t2};
 }
 
 // CHECK-LABEL: @main
 int main(int argc, char **argv) {
-// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{.+}}.region_id, i32 0, i8** null, i8** null, i64* null, i64* null)
+// CHECK: call i32 @__tgt_target_mapper(i64 -1, i8* @{{.+}}.region_id, i32 0, i8** null, i8** null, i64* null, i64* null, i8** null)
 #pragma omp target
  {
     for (int i = 0; i < 64; ++i) {
       for (int j = 0; j < 64; ++j) {
         auto foo = make_pair(i * i, j * j);
       }
     }
   }
   return 0;
 }
 
 // CHECK: call {{.+}} @{{.*}}make_pair

test/OpenMP/declare_mapper_codegen.cpp

 // RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK0 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 // RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 
 #ifdef CK0
 
 // CK0-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}.region_id = weak constant i8 0
-// CK0-64: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
-// CK0-32: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
+// CK0: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
 // CK0: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35]
-// CK0-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
-// CK0-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
+// CK0: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
 // CK0: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
-// CK0-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
-// CK0-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
+// CK0: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
 // CK0: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
 
 class C {
 public:
   int a;
   double *b;
 };
 
 #pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
 
-// CK0-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C.id{{.*}}(i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
+// CK0: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
 // CK0: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
 // CK0: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
 // CK0: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
 // CK0: store i64 %{{[^,]+}}, i{{64|32}}* [[SIZEADDR:%[^,]+]]
 // CK0: store i64 %{{[^,]+}}, i64* [[TYPEADDR:%[^,]+]]
 // CK0-DAG: [[SIZE:%.+]] = load i64, i64* [[SIZEADDR]]
 // CK0-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
 // CK0-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]]

 
 
 // CK0-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
 void foo(int a){
   int i = a;
   C c;
   c.a = a;
 
-  // CK0-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
+  // CK0-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
   // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
   // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
+  // CK0-DAG: [[MPRGEP]] = getelementptr inbounds {{.+}}[[MPR:%[^,]+]], i32 0, i32 0
   // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
   // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
+  // CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i32 0, i32 0
   // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
   // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
+  // CK0-DAG: [[CMPR1:%.+]] = bitcast i8** [[MPR1]] to void (i8*, i8*, i8*, i64, i64)**
   // CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
+  // CK0-DAG: store void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]], void (i8*, i8*, i8*, i64, i64)** [[CMPR1]]
   // CK0: call void [[KERNEL:@.+]](%class.C* [[VAL]])
   #pragma omp target map(mapper(id),tofrom: c)
   {
     ++c.a;
   }
 
-  // CK0-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}})
+  // CK0-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}}, i8** [[TMPRGEP:%.+]])
   // CK0-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
   // CK0-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+  // CK0-DAG: [[TMPRGEP]] = getelementptr inbounds {{.+}}[[TMPR:%[^,]+]], i32 0, i32 0
   // CK0-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0
   // CK0-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0
+  // CK0-DAG: [[TMPR1:%.+]] = getelementptr inbounds {{.+}}[[TMPR]], i32 0, i32 0
   // CK0-DAG: [[TCBP0:%.+]] = bitcast i8** [[TBP0]] to %class.C**
   // CK0-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C**
+  // CK0-DAG: [[TCMPR1:%.+]] = bitcast i8** [[TMPR1]] to void (i8*, i8*, i8*, i64, i64)**
   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCBP0]]
   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]]
+  // CK0-DAG: store void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]], void (i8*, i8*, i8*, i64, i64)** [[TCMPR1]]
   #pragma omp target update to(mapper(id): c)
 
-  // CK0-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}})
+  // CK0-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}}, i8** [[FMPRGEP:%.+]])
   // CK0-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
   // CK0-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+  // CK0-DAG: [[FMPRGEP]] = getelementptr inbounds {{.+}}[[FMPR:%[^,]+]], i32 0, i32 0
   // CK0-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0
   // CK0-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0
+  // CK0-DAG: [[FMPR1:%.+]] = getelementptr inbounds {{.+}}[[FMPR]], i32 0, i32 0
   // CK0-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C**
   // CK0-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C**
+  // CK0-DAG: [[FCMPR1:%.+]] = bitcast i8** [[FMPR1]] to void (i8*, i8*, i8*, i64, i64)**
   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]]
   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]]
+  // CK0-DAG: store void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]], void (i8*, i8*, i8*, i64, i64)** [[FCMPR1]]
   #pragma omp target update from(mapper(id): c)
 }
 
 
 // CK0: define internal void [[KERNEL]](%class.C* {{.+}}[[ARG:%.+]])
 // CK0: [[ADDR:%.+]] = alloca %class.C*,
 // CK0: store %class.C* [[ARG]], %class.C** [[ADDR]]
 // CK0: [[CADDR:%.+]] = load %class.C*, %class.C** [[ADDR]]
 // CK0: [[CAADDR:%.+]] = getelementptr inbounds %class.C, %class.C* [[CADDR]], i32 0, i32 0
 // CK0: [[VAL:%[^,]+]] = load i32, i32* [[CAADDR]]
 // CK0: {{.+}} = add nsw i32 [[VAL]], 1
 // CK0: }
 
-#endif
+#endif // CK0
 
 
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK1 --check-prefix CK1-64 %s
 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK1 --check-prefix CK1-64 %s
 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK1 --check-prefix CK1-32 %s
 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK1 --check-prefix CK1-32 %s
 
 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 
 #ifdef CK1
+// C++ template
 
 template <class T>
 class C {
 public:
   T a;
 };
 
 #pragma omp declare mapper(id: C<int> s) map(s.a)

 // CK1: [[DEL]]
 // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
 // CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
 // CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]])
 // CK1: br label %[[DONE]]
 // CK1: [[DONE]]
 // CK1: ret void
 
-#endif
+#endif // CK1
 
-#endif
+
+///==========================================================================///
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK2 --check-prefix CK2-64 %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK2 --check-prefix CK2-64 %s
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK2 --check-prefix CK2-32 %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK2 --check-prefix CK2-32 %s
+
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+
+#ifdef CK2
+// Nested mappers.
+
+class B {
+public:
+  double a;
+};
+
+class C {
+public:
+  double a;
+  B b;
+};
+
+#pragma omp declare mapper(B s) map(s.a)
+
+#pragma omp declare mapper(id: C s) map(s.b)
+
+// CK2: define {{.*}}void [[BMPRFUNC:@[.]omp_mapper[.].*B[.]default]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
+
+// CK2-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C{{.*}}.id(i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
+// CK2: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
+// CK2: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
+// CK2: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
+// CK2: store i64 %{{[^,]+}}, i{{64|32}}* [[SIZEADDR:%[^,]+]]
+// CK2: store i64 %{{[^,]+}}, i64* [[TYPEADDR:%[^,]+]]
+// CK2-DAG: [[SIZE:%.+]] = load i64, i64* [[SIZEADDR]]
+// CK2-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
+// CK2-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]]
+// CK2-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
+// CK2-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]]
+// CK2-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]]
+// CK2-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]]
+// CK2: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
+// CK2: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
+
+// CK2: [[INITEVALDEL]]
+// CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK2: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
+// CK2: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
+// CK2: [[INIT]]
+// CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
+// CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]])
+// CK2: br label %[[LHEAD:[^,]+]]
+
+// CK2: [[LHEAD]]
+// CK2: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
+// CK2: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
+// CK2: [[LBODY]]
+// CK2: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
+// CK2: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
+// CK2-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
+// CK2-DAG: [[BEND:%.+]] = getelementptr %class.B, %class.B* [[BBEGIN]], i32 1
+// CK2-DAG: [[BBEGINV:%.+]] = bitcast %class.B* [[BBEGIN]] to i8*
+// CK2-DAG: [[BENDV:%.+]] = bitcast %class.B* [[BEND]] to i8*
+// CK2-DAG: [[BBEGINI:%.+]] = ptrtoint i8* [[BBEGINV]] to i64
+// CK2-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64
+// CK2-DAG: [[BSIZE:%.+]] = sub i64 [[BENDI]], [[BBEGINI]]
+// CK2-DAG: [[BUSIZE:%.+]] = sdiv exact i64 [[BSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK2-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
+// CK2-DAG: [[PTRADDR0BC:%.+]] = bitcast %class.B* [[BBEGIN]] to i8*
+// CK2-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]])
+// CK2-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
+// CK2-DAG: br label %[[MEMBER:[^,]+]]
+// CK2-DAG: [[MEMBER]]
+// CK2-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
+// CK2-DAG: [[MEMBERCOM]]
+// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
+// CK2-DAG: br label %[[LTYPE]]
+// CK2-DAG: [[LTYPE]]
+// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK2-DAG: [[ALLOC]]
+// CK2-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK2-DAG: br label %[[TYEND:[^,]+]]
+// CK2-DAG: [[ALLOCELSE]]
+// CK2-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK2-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK2-DAG: [[TO]]
+// CK2-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
+// CK2-DAG: br label %[[TYEND]]
+// CK2-DAG: [[TOELSE]]
+// CK2-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK2-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK2-DAG: [[FROM]]
+// CK2-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
+// CK2-DAG: br label %[[TYEND]]
+// CK2-DAG: [[TYEND]]
+// CK2-DAG: [[TYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
+// CK2-64: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[BUSIZE]], i64 [[TYPE0]])
+// CK2-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
+// CK2-DAG: [[PTRADDR1BC:%.+]] = bitcast %class.B* [[BBEGIN]] to i8*
+// CK2-DAG: br label %[[MEMBER:[^,]+]]
+// CK2-DAG: [[MEMBER]]
+// CK2-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
+// CK2-DAG: [[MEMBERCOM]]
+// 281474976710659 == 0x1,000,000,003
+// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
+// CK2-DAG: br label %[[LTYPE]]
+// CK2-DAG: [[LTYPE]]
+// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK2-DAG: [[ALLOC]]
+// CK2-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK2-DAG: br label %[[TYEND:[^,]+]]
+// CK2-DAG: [[ALLOCELSE]]
+// CK2-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK2-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK2-DAG: [[TO]]
+// CK2-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
+// CK2-DAG: br label %[[TYEND]]
+// CK2-DAG: [[TOELSE]]
+// CK2-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK2-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK2-DAG: [[FROM]]
+// CK2-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
+// CK2-DAG: br label %[[TYEND]]
+// CK2-DAG: [[TYEND]]
+// CK2-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
+// CK2: call void [[BMPRFUNC]](i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 1, i64 [[TYPE1]])
+// CK2: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
+// CK2: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
+// CK2: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
+
+// CK2: [[LEXIT]]
+// CK2: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
+// CK2: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
+// CK2: [[EVALDEL]]
+// CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK2: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
+// CK2: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
+// CK2: [[DEL]]
+// CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
+// CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]])
+// CK2: br label %[[DONE]]
+// CK2: [[DONE]]
+// CK2: ret void
+
+#endif // CK2
+
+#endif // HEADER

test/OpenMP/declare_target_link_codegen.cpp

 // HOST: [[BP2_CAST:%.+]] = bitcast i8** [[BP2]] to i32***
 // HOST: store i32** @[[D_PTR]], i32*** [[BP2_CAST]],
 // HOST: [[P2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
 // HOST: [[P2_CAST:%.+]] = bitcast i8** [[P2]] to i32**
 // HOST: store i32* @[[D]], i32** [[P2_CAST]],
 
 // HOST: [[BP0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
 // HOST: [[P0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
-// HOST: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP0]], i8** [[P0]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[SIZES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPTYPES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0))
+// HOST: call i32 @__tgt_target_mapper(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP0]], i8** [[P0]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[SIZES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPTYPES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0))
 // HOST: call void @__omp_offloading_{{.*}}_{{.*}}_{{.*}}maini1{{.*}}_l42(i32* %{{[^,]+}})
-// HOST: call i32 @__tgt_target_teams(i64 -1, i8* @.__omp_offloading_{{.+}}_l47.region_id, i32 2, {{.+}})
+// HOST: call i32 @__tgt_target_teams_mapper(i64 -1, i8* @.__omp_offloading_{{.+}}_l47.region_id, i32 2, {{.+}})
 
 // HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l42(i32* dereferenceable{{.*}})
 // HOST: [[C:%.*]] = load i32, i32* @c,
 // HOST: store i32 [[C]], i32* %
 
 // CHECK: !{i32 1, !"c_decl_tgt_ref_ptr", i32 1, i32 {{[0-9]+}}}
 #endif // HEADER

test/OpenMP/target_is_device_ptr_codegen.cpp

 // CK1: [[TYPES06:@.+]] = {{.+}}constant [2 x i64] [i64 288, i64 288]
 
 // CK1-LABEL: @_Z3foo{{.*}}(
 template<typename T>
 void foo(float *&lr, T *&tr) {
   float *l;
   T *t;
 
-  // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES00]]{{.+}}, {{.+}}[[TYPES00]]{{.+}})
+  // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES00]]{{.+}}, {{.+}}[[TYPES00]]{{.+}})
   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
   // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double**
   // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double**
   // CK1-DAG: store double* [[VAL:%.+]], double** [[CBP1]]
   // CK1-DAG: store double* [[VAL]], double** [[CP1]]
   // CK1-DAG: [[VAL]] = load double*, double** [[ADDR:@g]],
 
   // CK1: call void [[KERNEL:@.+]](double* [[VAL]])
   #pragma omp target is_device_ptr(g)
   {
     ++g;
   }
 
-  // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES01]]{{.+}}, {{.+}}[[TYPES01]]{{.+}})
+  // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES01]]{{.+}}, {{.+}}[[TYPES01]]{{.+}})
   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
   // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to float**
   // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to float**
   // CK1-DAG: store float* [[VAL:%.+]], float** [[CBP1]]
   // CK1-DAG: store float* [[VAL]], float** [[CP1]]
   // CK1-DAG: [[VAL]] = load float*, float** [[ADDR:%.+]],
 
   // CK1: call void [[KERNEL:@.+]](float* [[VAL]])
   #pragma omp target is_device_ptr(l)
   {
     ++l;
   }
 
-  // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES02]]{{.+}}, {{.+}}[[TYPES02]]{{.+}})
+  // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES02]]{{.+}}, {{.+}}[[TYPES02]]{{.+}})
   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
   // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32**
   // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
   // CK1-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]]
   // CK1-DAG: store i32* [[VAL]], i32** [[CP1]]
   // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
 
   // CK1: call void [[KERNEL:@.+]](i32* [[VAL]])
   #pragma omp target is_device_ptr(t)
   {
     ++t;
   }
 
-  // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES03]]{{.+}}, {{.+}}[[TYPES03]]{{.+}})
+  // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES03]]{{.+}}, {{.+}}[[TYPES03]]{{.+}})
   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
   // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to float**
   // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to float**
   // CK1-DAG: store float* [[VAL:%.+]], float** [[CBP1]]
   // CK1-DAG: store float* [[VAL]], float** [[CP1]]
   // CK1-DAG: [[VAL]] = load float*, float** [[ADDR:%.+]],
   // CK1-DAG: [[ADDR]] = load float**, float*** [[ADDR2:%.+]],
 
   // CK1: call void [[KERNEL:@.+]](float* [[VAL]])
   #pragma omp target is_device_ptr(lr)
   {
     ++lr;
   }
 
-  // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES04]]{{.+}}, {{.+}}[[TYPES04]]{{.+}})
+  // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES04]]{{.+}}, {{.+}}[[TYPES04]]{{.+}})
   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
   // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32**
   // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
   // CK1-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]]
   // CK1-DAG: store i32* [[VAL]], i32** [[CP1]]
   // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
   // CK1-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
 
   // CK1: call void [[KERNEL:@.+]](i32* [[VAL]])
   #pragma omp target is_device_ptr(tr)
   {
     ++tr;
   }
 
-  // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES05]]{{.+}}, {{.+}}[[TYPES05]]{{.+}})
+  // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES05]]{{.+}}, {{.+}}[[TYPES05]]{{.+}})
   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
   // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32**
   // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
   // CK1-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]]
   // CK1-DAG: store i32* [[VAL]], i32** [[CP1]]
   // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
   // CK1-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
 
   // CK1: call void [[KERNEL:@.+]](i32* [[VAL]])
   #pragma omp target is_device_ptr(tr,lr)
   {
     ++tr;
   }
 
-  // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES06]]{{.+}}, {{.+}}[[TYPES06]]{{.+}})
+  // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES06]]{{.+}}, {{.+}}[[TYPES06]]{{.+}})
   // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
   // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
   // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32**
   // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
   // CK1-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]]
   // CK1-DAG: store i32* [[VAL]], i32** [[CP1]]
struct ST {
   T *a;
   double *&b;
   ST(double *&b) : a(0), b(b) {}
 
   // CK2-LABEL: @{{.*}}foo{{.*}}
   void foo(double *&arg) {
     int *la = 0;
 
-    // CK2-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+    // CK2-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
     // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
     // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
     // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
     // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
     // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
     // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
     // CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]]
     // CK2-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CP0]]
     #pragma omp target is_device_ptr(a)
     {
       a++;
     }
 
-    // CK2-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+    // CK2-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
     // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
     // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
     // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
     // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
     // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
     // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
     // CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]]
     // CK2-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CP0]]
     #pragma omp target is_device_ptr(b)
     {
       b++;
     }
 
-    // CK2-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
+    // CK2-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
     // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
     // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
     // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
     // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
     // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
     // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
     // CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]]