[OpenMPOpt][SplitMemTransfer][WIP] Splitting the runtime call __tgt_target_data_begin

[hamax97]

Summary

Finally the functionality to split the runtimecall __tgt_target_data_begin.

• Still problems querying AAResults.

[jdoerfert]

This still conflates different tasks and should be separated. The definitions of the new issue and wait function need to be provided with the splitting code. Just splitting, not all the other stuff. We can then also test that.

llvm/lib/Transforms/IPO/OpenMPOpt.cpp

678–679

Also below, no need to check first.

1045

Don't mutate, I have never seen this actually. Create a new call and remove the old one.

1051

You don't need a dynamic array here and below.

1084–1086

1097

If you define Before = After->getNextNode(); you can insertBefore(Before) without resetting the insertion point all the time. Isn't Issues unordered?

[hamax97]

I think if I only add the declarations of __tgt_target_data_begin_issue and __tgt_target_data_being_wait the linking process will fail. How should I define the wait function?, the issue is just a wrapper of the current function, but as we don't have Shilei's patch merged yet I would have to create something that does nothing and change it later when the asynchronous stuff is merged.

llvm/lib/Transforms/IPO/OpenMPOpt.cpp

678–679

Sure, I'll change it, I forgot SetVector ignore repeated elements.

1045

Ok, I'll change that.

1051

I didn't know what to do. I think we need a long life array because the array will be erased when the function scope ends, won't it?.

1084–1086

I'll fix it.

1097

Ohh sure, that is handier. But Issue is not unordered, if we mess up the order of the instructions we change the semantics. What is nice about SetVector is that it allows you to visit its elements in the order they were inserted.

[jdoerfert]

In D85535#2207439, @hamax97 wrote:

I think if I only add the declarations of __tgt_target_data_begin_issue and __tgt_target_data_being_wait the linking process will fail. How should I define the wait function?, the issue is just a wrapper of the current function, but as we don't have Shilei's patch merged yet I would have to create something that does nothing and change it later when the asynchronous stuff is merged.

We can add this under a command line flag, declarations will be sufficient to test it with lit tests then.

[jdoerfert]

llvm/lib/Transforms/IPO/OpenMPOpt.cpp

1051

It will but the call copies the entries, that is why it takes an ArrayRef, bascially something that doesn't promise a lifetime. If you need the content, copy it.

1097

Ordered makes sense ;)

[hamax97]

• Using OMPIRBuilder instead of manually generating the runtime calls.
• Updating regression test for this optimization.

[hamax97]

• Splitting patch into canBeMoved logic and splitting the runtime call.
• Fix to minor bug in detectIssue.
• Update regression test. Now checks that the runtime call __tgt_target_data_begin is successfully split. Not sure if this should go in another patch.

[jdoerfert]

Is this a diff against master?

llvm/lib/Transforms/IPO/OpenMPOpt.cpp

1072

Let's remove all the move code from this one.

[hamax97]

In D85535#2213345, @jdoerfert wrote:

Is this a diff against master?

No, it's against D84593. And that one is against D82719.

[hamax97]

• Removing move code.

[jdoerfert]

Hm, is it possible to apply this one first. I think this is the easiest to approve right away, that way we have less outstanding code.

[sstefan1]

llvm/lib/Transforms/IPO/OpenMPOpt.cpp

1039

Maybe SmallVector?

[hamax97]

• Changing std::vector for SmallVector in split function.

[jdoerfert]

In D85535#2213698, @jdoerfert wrote:

Hm, is it possible to apply this one first. I think this is the easiest to approve right away, that way we have less outstanding code.

WDYT?

[hamax97]

In D85535#2214821, @jdoerfert wrote:

In D85535#2213698, @jdoerfert wrote:

Hm, is it possible to apply this one first. I think this is the easiest to approve right away, that way we have less outstanding code.

WDYT?

But we need to add the structure MemoryTranfer too. Also remember this assumes we moved the declarations in OpenMPOpt.cpp to OpenMPOpt.h, so we would have to put this new code into OpenMPOpt.cpp, which I think is not a problem at all, but just to consider it.

[jdoerfert]

In D85535#2215712, @hamax97 wrote:

In D85535#2214821, @jdoerfert wrote:

In D85535#2213698, @jdoerfert wrote:

Hm, is it possible to apply this one first. I think this is the easiest to approve right away, that way we have less outstanding code.

WDYT?

But we need to add the structure MemoryTranfer too. Also remember this assumes we moved the declarations in OpenMPOpt.cpp to OpenMPOpt.h, so we would have to put this new code into OpenMPOpt.cpp, which I think is not a problem at all, but just to consider it.

Why do we need anything to split the runtime calls? The arguments should be sufficient for the split, no?

[hamax97]

In D85535#2215750, @jdoerfert wrote:

In D85535#2215712, @hamax97 wrote:

In D85535#2214821, @jdoerfert wrote:

In D85535#2213698, @jdoerfert wrote:

Hm, is it possible to apply this one first. I think this is the easiest to approve right away, that way we have less outstanding code.

WDYT?

But we need to add the structure MemoryTranfer too. Also remember this assumes we moved the declarations in OpenMPOpt.cpp to OpenMPOpt.h, so we would have to put this new code into OpenMPOpt.cpp, which I think is not a problem at all, but just to consider it.

Why do we need anything to split the runtime calls? The arguments should be sufficient for the split, no?

In that case, sure. I will do that.

[hamax97]

• Isolating split function.

[jdoerfert]

llvm/lib/Transforms/IPO/OpenMPOpt.cpp

48

Call the variable and option "hide memory transfer latency" (or similar) and add a WIP to the description.

692

This is not the correct return value. The lambda needs to return if the use was deleted or not. So you need to track "globally changed" (which is Changed right now), and "use changed" (which you return from the lambda).

700

We need a more descriptive name here, e.g., splitTargetDataBeginRTC.

710

Let's do 8 not 7, no need for the reserve.

715

Nit: No need for the ArrayRef.

716–717

724–731

llvm/test/Transforms/OpenMP/hide_mem_transfer_latency.ll

2

Can you make an NFC pre-commit that updates the test, and then only include the -openmp-split-memtransfers changes in this patch.

[hamax97]

llvm/lib/Transforms/IPO/OpenMPOpt.cpp

48

Sure

692

Sure, I'll fix it. If I understood I need to return from the lambda the immediate value of split, without forgetting to make an or with the outer Changed.

700

Sure.

llvm/test/Transforms/OpenMP/hide_mem_transfer_latency.ll

2

Sure, but the test will fail, we would have to commit them at the same time.

[hamax97]

• Refactor to some names and other minor changes.
• Update regression test to expect the split runtime calls.

[jdoerfert]

LGTM, please clang format the code.

[jdoerfert]

update the commit message before merging this

[hamax97]

Commit: https://github.com/llvm/llvm-project/commit/496f8e5b369f091def93482578232da8c6e77a7a

llvm/include/llvm/Frontend/OpenMP/OMPKinds.def

 #ifndef OMP_STRUCT_TYPE
 #define OMP_STRUCT_TYPE(VarName, StructName, ...)
 #endif
 
 #define __OMP_STRUCT_TYPE(VarName, Name, ...)                                  \
   OMP_STRUCT_TYPE(VarName, "struct." #Name, __VA_ARGS__)
 
 __OMP_STRUCT_TYPE(Ident, ident_t, Int32, Int32, Int32, Int32, Int8Ptr)
+__OMP_STRUCT_TYPE(AsyncInfo, __tgt_async_info, Int8Ptr)
 
 #undef __OMP_STRUCT_TYPE
 #undef OMP_STRUCT_TYPE
 
 #ifndef OMP_FUNCTION_TYPE
 #define OMP_FUNCTION_TYPE(VarName, IsVarArg, ReturnType, ...)
 #endif
 
__OMP_RTL(__tgt_target_teams_mapper, false, Int32, Int64, VoidPtr, Int32,
           VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, Int32, Int32)
 __OMP_RTL(__tgt_target_teams_nowait_mapper, false, Int32, Int64, VoidPtr, Int32,
           VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr, Int32, Int32)
 __OMP_RTL(__tgt_register_requires, false, Void, Int64)
 __OMP_RTL(__tgt_target_data_begin_mapper, false, Void, Int64, Int32, VoidPtrPtr,
           VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
 __OMP_RTL(__tgt_target_data_begin_nowait_mapper, false, Void, Int64, Int32,
           VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
+__OMP_RTL(__tgt_target_data_begin_mapper_issue, false, AsyncInfo, Int64, Int32,
+          VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
+__OMP_RTL(__tgt_target_data_begin_mapper_wait, false, Void, Int64, AsyncInfo)
 __OMP_RTL(__tgt_target_data_end_mapper, false, Void, Int64, Int32, VoidPtrPtr,
           VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
 __OMP_RTL(__tgt_target_data_end_nowait_mapper, false, Void, Int64, Int32,
           VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
 __OMP_RTL(__tgt_target_data_update_mapper, false, Void, Int64, Int32,
           VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
 __OMP_RTL(__tgt_target_data_update_nowait_mapper, false, Void, Int64, Int32,
           VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)

llvm/lib/Transforms/IPO/OpenMPOpt.cpp

static cl::opt<bool> DisableOpenMPOptimizations(
     cl::desc("Disable OpenMP specific optimizations."), cl::Hidden,
     cl::init(false));
 
 static cl::opt<bool> PrintICVValues("openmp-print-icv-values", cl::init(false),
                                     cl::Hidden);
 static cl::opt<bool> PrintOpenMPKernels("openmp-print-gpu-kernels",
                                         cl::init(false), cl::Hidden);
 
+static cl::opt<bool> SplitMemoryTransfers(
+    "openmp-split-memtransfers",
+    cl::desc("Tries to hide the latency of host to device memory transfers"),
+    cl::Hidden, cl::init(false));
+
+
 STATISTIC(NumOpenMPRuntimeCallsDeduplicated,
           "Number of OpenMP runtime calls deduplicated");
 STATISTIC(NumOpenMPParallelRegionsDeleted,
           "Number of OpenMP parallel regions deleted");
 STATISTIC(NumOpenMPRuntimeFunctionsIdentified,
           "Number of OpenMP runtime functions identified");
 STATISTIC(NumOpenMPRuntimeFunctionUsesIdentified,
           "Number of OpenMP runtime function uses identified");
bool run() {
 
     Changed |= runAttributor();
 
     // Recollect uses, in case Attributor deleted any.
     OMPInfoCache.recollectUses();
 
     Changed |= deduplicateRuntimeCalls();
     Changed |= deleteParallelRegions();
+    if (SplitMemoryTransfers)
+      Changed |= hideMemTransfersLatency();
 
     return Changed;
   }
 
   /// Print initial ICV values for testing.
   /// FIXME: This should be done from the Attributor once it is added.
   void printICVs() const {
     InternalControlVar ICVs[] = {ICV_nthreads, ICV_active_levels, ICV_cancel};
for (Function *F : SCC) {
         }
       Changed |= deduplicateRuntimeCalls(
           *F, OMPInfoCache.RFIs[OMPRTL___kmpc_global_thread_num], GTIdArg);
     }
 
     return Changed;
   }
 
+  /// Tries to hide the latency of runtime calls that involve host to
+  /// device memory transfers by splitting them into their "issue" and "wait"
+  /// versions. The "issue" is moved upwards as much as possible. The "wait" is
+  /// moved downards as much as possible. The "issue" issues the memory transfer
+  /// asynchronously, returning a handle. The "wait" waits in the returned
+  /// handle for the memory transfer to finish.
+  bool hideMemTransfersLatency() {
+    auto &RFI = OMPInfoCache.RFIs[OMPRTL___tgt_target_data_begin_mapper];
+    bool Changed = false;
+    auto SplitMemTransfers = [&](Use &U, Function &Decl) {
+      auto *RTCall = getCallIfRegularCall(U, &RFI);
+      if (!RTCall)
+        return false;
+
+      Changed |= split(RTCall);
+      return Changed;
+    };
+    RFI.foreachUse(SCC, SplitMemTransfers);
+
+    return Changed;
+  }
+
+  /// Splits \p RuntimeCall into its "issue" and "wait" counterparts.
+  bool split(CallInst *RuntimeCall) {
+    auto &IRBuilder = OMPInfoCache.OMPBuilder;
+    // Add "issue" runtime call declaration.
+    // declare %struct.tgt_async_info @__tgt_target_data_begin_issue(i64, i32,
+    //   i8**, i8**, i64*, i64*)
+    FunctionCallee IssueDecl = IRBuilder.getOrCreateRuntimeFunction(
+        M, OMPRTL___tgt_target_data_begin_mapper_issue);
+
+    // Change RuntimeCall callsite for its asynchronous version.
+    SmallVector<Value *, 7> Args;
+    Args.reserve(RuntimeCall->getNumArgOperands());
+    for (auto &Arg : RuntimeCall->args())
+      Args.push_back(Arg.get());
+
+    CallInst *IssueCallsite = CallInst::Create(
+        IssueDecl, ArrayRef<Value *>(Args), "handle", RuntimeCall);
+    RuntimeCall->removeFromParent();
+    RuntimeCall->deleteValue();
+
+    // Add "wait" runtime call declaration.
+    // declare void @__tgt_target_data_begin_wait(i64, %struct.__tgt_async_info)
+    FunctionCallee WaitDecl = IRBuilder.getOrCreateRuntimeFunction(
+        M, OMPRTL___tgt_target_data_begin_mapper_wait);
+
+    // Add "wait" call site.
+    const unsigned WaitNumParams = 2;
+    Value *WaitParams[] = {
+        IssueCallsite->getArgOperand(0), // device_id.
+        IssueCallsite // returned handle.
+    };
+    CallInst::Create(
+        WaitDecl, ArrayRef<Value *>(WaitParams, WaitNumParams), /*NameStr=*/"",
+            IssueCallsite->getNextNode());
+
+    return true;
+  }
+
   static Value *combinedIdentStruct(Value *CurrentIdent, Value *NextIdent,
                                     bool GlobalOnly, bool &SingleChoice) {
     if (CurrentIdent == NextIdent)
       return CurrentIdent;
 
     // TODO: Figure out how to actually combine multiple debug locations. For
     //       now we just keep an existing one if there is a single choice.
     if (!GlobalOnly || isa<GlobalValue>(NextIdent)) {

 Kernel OpenMPOpt::getUniqueKernelFor(Function &F) {
   if (!OMPInfoCache.ModuleSlice.count(&F))
     return nullptr;
 
   // Use a scope to keep the lifetime of the CachedKernel short.
   {
     Optional<Kernel> &CachedKernel = UniqueKernelMap[&F];
     if (CachedKernel)
       return *CachedKernel;
 
     // TODO: We should use an AA to create an (optimistic and callback
     //       call-aware) call graph. For now we stick to simple patterns that
     //       are less powerful, basically the worst fixpoint.
     if (isKernel(F)) {
       CachedKernel = Kernel(&F);
       return *CachedKernel;
     }
 
     CachedKernel = nullptr;
     if (!F.hasLocalLinkage())
       return nullptr;
   }
 
   auto GetUniqueKernelForUse = [&](const Use &U) -> Kernel {
     if (auto *Cmp = dyn_cast<ICmpInst>(U.getUser())) {
       // Allow use in equality comparisons.
       if (Cmp->isEquality())
         return getUniqueKernelFor(*Cmp);
       return nullptr;
     }
     if (auto *CB = dyn_cast<CallBase>(U.getUser())) {
       // Allow direct calls.
       if (CB->isCallee(&U))
         return getUniqueKernelFor(*CB);
       // Allow the use in __kmpc_kernel_prepare_parallel calls.
       if (Function *Callee = CB->getCalledFunction())
         if (Callee->getName() == "__kmpc_kernel_prepare_parallel")
           return getUniqueKernelFor(*CB);
       return nullptr;
     }
     // Disallow every other use.
     return nullptr;
   };
 
   // TODO: In the future we want to track more than just a unique kernel.
   SmallPtrSet<Kernel, 2> PotentialKernels;
   foreachUse(F, [&](const Use &U) {
     PotentialKernels.insert(GetUniqueKernelForUse(U));
   });
 
   Kernel K = nullptr;
   if (PotentialKernels.size() == 1)
     K = *PotentialKernels.begin();
 
   // Cache the result.
   UniqueKernelMap[&F] = K;
 
   return K;
 }
 
 bool OpenMPOpt::rewriteDeviceCodeStateMachine() {
   OMPInformationCache::RuntimeFunctionInfo &KernelPrepareParallelRFI =
       OMPInfoCache.RFIs[OMPRTL___kmpc_kernel_prepare_parallel];
 
   bool Changed = false;
   if (!KernelPrepareParallelRFI)
     return Changed;
 
   for (Function *F : SCC) {
 
     // Check if the function is uses in a __kmpc_kernel_prepare_parallel call at
     // all.
     bool UnknownUse = false;
     bool KernelPrepareUse = false;
     unsigned NumDirectCalls = 0;

llvm/test/Transforms/OpenMP/hide_mem_transfer_latency.ll

-; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: -p --function-signature
-; RUN: opt -S -passes=openmpopt < %s | FileCheck %s
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: -p --function-signature --scrub-attributes
+; RUN: opt -S -passes=openmpopt -aa-pipeline=basic-aa -openmp-split-memtransfers < %s | FileCheck %s
 target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
 
-; FIXME: This struct should be generated after splitting at least one of the runtime calls.
-; %struct.__tgt_async_info = type { i8* }
+; CHECK: %struct.__tgt_async_info = type { i8* }
+
 %struct.ident_t = type { i32, i32, i32, i32, i8* }
 %struct.__tgt_offload_entry = type { i8*, i8*, i64, i32, i32 }
 
 @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 35]
-@.__omp_offloading_heavyComputation1.region_id = weak constant i8 0
+@.__omp_offloading_heavyComputation1_region_id = weak constant i8 0
 @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 8]
 @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 800]
-
-@.__omp_offloading_heavyComputation2.region_id = weak constant i8 0
-@.offload_maptypes.3 = private unnamed_addr constant [2 x i64] [i64 35, i64 35]
-
-@.__omp_offloading_heavyComputation3.region_id = weak constant i8 0
-@.offload_sizes.2 = private unnamed_addr constant [2 x i64] [i64 4, i64 0]
-@.offload_maptypes.4 = private unnamed_addr constant [2 x i64] [i64 800, i64 544]
-
-@.offload_maptypes.5 = private unnamed_addr constant [1 x i64] [i64 33]
+@.__omp_offloading_heavyComputation2_region_id = weak constant i8 0
+@.offload_maptypes.7 = private unnamed_addr constant [2 x i64] [i64 35, i64 35]
+@.__omp_offloading_heavyComputation3_region_id = weak constant i8 0
+@.offload_sizes.9 = private unnamed_addr constant [2 x i64] [i64 4, i64 0]
+@.offload_maptypes.10 = private unnamed_addr constant [2 x i64] [i64 800, i64 544]
+@.offload_maptypes.11 = private unnamed_addr constant [1 x i64] [i64 33]
 
 ;double heavyComputation1() {
 ;  double a = rand() % 777;
 ;  double random = rand();
 ;
 ;  //#pragma omp target data map(a)
 ;  void* args[1];
 ;  args[0] = &a;

 ; CHECK-LABEL: define {{[^@]+}}@heavyComputation1()
 ; CHECK-NEXT:  entry:
 ; CHECK-NEXT:    %a = alloca double, align 8
 ; CHECK-NEXT:    %.offload_baseptrs = alloca [1 x i8*], align 8
 ; CHECK-NEXT:    %.offload_ptrs = alloca [1 x i8*], align 8
 ; CHECK-NEXT:    %.offload_baseptrs4 = alloca [1 x i8*], align 8
 ; CHECK-NEXT:    %.offload_ptrs5 = alloca [1 x i8*], align 8
 ; CHECK-NEXT:    %0 = bitcast double* %a to i8*
-; CHECK-NEXT:    %call = call i32 @rand()
+; CHECK-NEXT:    %call = tail call i32 (...) @rand()
 ; CHECK-NEXT:    %rem = srem i32 %call, 777
 ; CHECK-NEXT:    %conv = sitofp i32 %rem to double
 ; CHECK-NEXT:    store double %conv, double* %a, align 8
-; CHECK-NEXT:    %call1 = call i32 @rand()
+; CHECK-NEXT:    %call1 = tail call i32 (...) @rand()
 ; CHECK-NEXT:    %1 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_baseptrs, i64 0, i64 0
 ; CHECK-NEXT:    %2 = bitcast [1 x i8*]* %.offload_baseptrs to double**
 ; CHECK-NEXT:    store double* %a, double** %2, align 8
 ; CHECK-NEXT:    %3 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i64 0, i64 0
 ; CHECK-NEXT:    %4 = bitcast [1 x i8*]* %.offload_ptrs to double**
 ; CHECK-NEXT:    store double* %a, double** %4, align 8
-; CHECK-NEXT:    call void @__tgt_target_data_begin(i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i64 0, i64 0))
+
+; CHECK-NEXT:    %handle = call %struct.__tgt_async_info @__tgt_target_data_begin_mapper_issue(i64 -1, i32 1, i8** %1, i8** %3, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i64 0, i64 0), i8** null)
+; CHECK-NEXT:    call void @__tgt_target_data_begin_mapper_wait(i64 -1, %struct.__tgt_async_info %handle)
+
 ; CHECK-NEXT:    %5 = bitcast double* %a to i64*
 ; CHECK-NEXT:    %6 = load i64, i64* %5, align 8
 ; CHECK-NEXT:    %7 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_baseptrs4, i64 0, i64 0
 ; CHECK-NEXT:    %8 = bitcast [1 x i8*]* %.offload_baseptrs4 to i64*
 ; CHECK-NEXT:    store i64 %6, i64* %8, align 8
 ; CHECK-NEXT:    %9 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs5, i64 0, i64 0
 ; CHECK-NEXT:    %10 = bitcast [1 x i8*]* %.offload_ptrs5 to i64*
 ; CHECK-NEXT:    store i64 %6, i64* %10, align 8
-; CHECK-NEXT:    %11 = call i32 @__tgt_target_teams(i64 -1, i8* nonnull @.__omp_offloading_heavyComputation1.region_id, i32 1, i8** nonnull %7, i8** nonnull %9, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.2, i64 0, i64 0), i32 0, i32 0)
-; CHECK-NEXT:    %12 = icmp eq i32 %11, 0
-; CHECK-NEXT:    br i1 %12, label %omp_offload.cont, label %omp_offload.failed
+; CHECK-NEXT:    %11 = call i32 @__tgt_target_teams_mapper(i64 -1, i8* nonnull @.__omp_offloading_heavyComputation1_region_id, i32 1, i8** nonnull %7, i8** nonnull %9, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.2, i64 0, i64 0), i8** null, i32 0, i32 0)
+; CHECK-NEXT:    %.not = icmp eq i32 %11, 0
+; CHECK-NEXT:    br i1 %.not, label %omp_offload.cont, label %omp_offload.failed
 ; CHECK:       omp_offload.failed:
 ; CHECK-NEXT:    call void @heavyComputation1FallBack(i64 %6)
 ; CHECK-NEXT:    br label %omp_offload.cont
 ; CHECK:       omp_offload.cont:
 ; CHECK-NEXT:    %conv2 = sitofp i32 %call1 to double
-; CHECK-NEXT:    call void @__tgt_target_data_end(i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i64 0, i64 0))
-; CHECK-NEXT:    %13 = load double, double* %a, align 8
-; CHECK-NEXT:    %add = fadd double %13, %conv2
+; CHECK-NEXT:    call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i64 0, i64 0), i8** null)
+; CHECK-NEXT:    %12 = load double, double* %a, align 8
+; CHECK-NEXT:    %add = fadd double %12, %conv2
 ; CHECK-NEXT:    ret double %add
 ;
 entry:
   %a = alloca double, align 8
   %.offload_baseptrs = alloca [1 x i8*], align 8
   %.offload_ptrs = alloca [1 x i8*], align 8
   %.offload_baseptrs4 = alloca [1 x i8*], align 8
   %.offload_ptrs5 = alloca [1 x i8*], align 8
 
-  ; FIXME: Should have after splitting the runtime call __tgt_target_data_begin.
-  ; %device_id1 = alloca i64, align 8
-  ; %async_info1 = alloca %struct.__tgt_async_info, align 8
-
   %0 = bitcast double* %a to i8*
-  %call = call i32 @rand()
+  %call = tail call i32 (...) @rand()
   %rem = srem i32 %call, 777
   %conv = sitofp i32 %rem to double
   store double %conv, double* %a, align 8
 
-  ; FIXME: The "isue" should be moved here.
-  %call1 = call i32 @rand()
+  ; FIXME: call i8* @__tgt_target_data_begin_mapper_issue(...) should be moved here.
+  %call1 = tail call i32 (...) @rand()
 
-  ; FIXME: This setup for the runtime call __tgt_target_data_begin should be
-  ;        split into its "issue" and "wait" counterpars and moved upwards
-  ;        and downwards, respectively. The call should be replaced to something
-  ;        like ...
-  ; Issue - this is moved upwards.
-  ; ... setup code ...
-  ; store i64 -1, i64* %device_id1, align 8
-  ; %handle1 = call i8* @__tgt_target_data_begin(i64* dereferenceable(8) %device_id1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i64 0, i64 0))
-  ; Wait - this is moved downwards.
-  ; %device_id1_copy = load i64, i64* %device_id1, align 8 ; device_id
-  ; %queue1 = getelementptr inbounds %struct.__tgt_async_info, %struct.__tgt_async_info* %async_info1, i32 0, i32 0
-  ; store i8* %handle1, i8** %queue1, align 8
-  ; call void @__tgt_target_data_begin_wait(i64 %device_id1_copy, %struct.__tgt_async_info* dereferenceable(8) %async_info1)
   %1 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_baseptrs, i64 0, i64 0
   %2 = bitcast [1 x i8*]* %.offload_baseptrs to double**
   store double* %a, double** %2, align 8
   %3 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i64 0, i64 0
   %4 = bitcast [1 x i8*]* %.offload_ptrs to double**
   store double* %a, double** %4, align 8
-  call void @__tgt_target_data_begin(i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i64 0, i64 0))
+  call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i64 0, i64 0), i8** null)
 
   %5 = bitcast double* %a to i64*
   %6 = load i64, i64* %5, align 8
   %7 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_baseptrs4, i64 0, i64 0
   %8 = bitcast [1 x i8*]* %.offload_baseptrs4 to i64*
   store i64 %6, i64* %8, align 8
   %9 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs5, i64 0, i64 0
   %10 = bitcast [1 x i8*]* %.offload_ptrs5 to i64*
   store i64 %6, i64* %10, align 8
 
-  ; FIXME: The "wait" should be moved here.
-  %11 = call i32 @__tgt_target_teams(i64 -1, i8* nonnull @.__omp_offloading_heavyComputation1.region_id, i32 1, i8** nonnull %7, i8** nonnull %9, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.2, i64 0, i64 0), i32 0, i32 0)
-
-  %12 = icmp eq i32 %11, 0
-  br i1 %12, label %omp_offload.cont, label %omp_offload.failed
+  ; FIXME: call i8* @__tgt_target_data_begin_mapper_wait(...) should be moved here.
+  %11 = call i32 @__tgt_target_teams_mapper(i64 -1, i8* nonnull @.__omp_offloading_heavyComputation1_region_id, i32 1, i8** nonnull %7, i8** nonnull %9, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.2, i64 0, i64 0), i8** null, i32 0, i32 0)
+  %.not = icmp eq i32 %11, 0
+  br i1 %.not, label %omp_offload.cont, label %omp_offload.failed
 
 omp_offload.failed:                               ; preds = %entry
   call void @heavyComputation1FallBack(i64 %6)
   br label %omp_offload.cont
 
-omp_offload.cont:                                 ; preds = %entry, %omp_offload.failed
+omp_offload.cont:                                 ; preds = %omp_offload.failed, %entry
   %conv2 = sitofp i32 %call1 to double
-  call void @__tgt_target_data_end(i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i64 0, i64 0))
-  %13 = load double, double* %a, align 8
-  %add = fadd double %13, %conv2
+  call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i64 0, i64 0), i8** null)
+  %12 = load double, double* %a, align 8
+  %add = fadd double %12, %conv2
   ret double %add
 }
 
 define internal void @heavyComputation1FallBack(i64 %a) {
 entry:
   ; Fallback for offloading function heavyComputation1.
   ret void
 }

 ; CHECK-NEXT:  entry:
 ; CHECK-NEXT:    %size.addr = alloca i32, align 4
 ; CHECK-NEXT:    %.offload_baseptrs = alloca [2 x i8*], align 8
 ; CHECK-NEXT:    %.offload_ptrs = alloca [2 x i8*], align 8
 ; CHECK-NEXT:    %.offload_sizes = alloca [2 x i64], align 8
 ; CHECK-NEXT:    %.offload_baseptrs2 = alloca [2 x i8*], align 8
 ; CHECK-NEXT:    %.offload_ptrs3 = alloca [2 x i8*], align 8
 ; CHECK-NEXT:    store i32 %size, i32* %size.addr, align 4
-; CHECK-NEXT:    %call = call i32 @rand()
+; CHECK-NEXT:    %call = tail call i32 (...) @rand()
 ; CHECK-NEXT:    %conv = zext i32 %size to i64
 ; CHECK-NEXT:    %0 = shl nuw nsw i64 %conv, 3
 ; CHECK-NEXT:    %1 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 0
 ; CHECK-NEXT:    %2 = bitcast [2 x i8*]* %.offload_baseptrs to double**
 ; CHECK-NEXT:    store double* %a, double** %2, align 8
 ; CHECK-NEXT:    %3 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 0
 ; CHECK-NEXT:    %4 = bitcast [2 x i8*]* %.offload_ptrs to double**
 ; CHECK-NEXT:    store double* %a, double** %4, align 8
 ; CHECK-NEXT:    %5 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 0
 ; CHECK-NEXT:    store i64 %0, i64* %5, align 8
 ; CHECK-NEXT:    %6 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 1
 ; CHECK-NEXT:    %7 = bitcast i8** %6 to i32**
 ; CHECK-NEXT:    store i32* %size.addr, i32** %7, align 8
 ; CHECK-NEXT:    %8 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 1
 ; CHECK-NEXT:    %9 = bitcast i8** %8 to i32**
 ; CHECK-NEXT:    store i32* %size.addr, i32** %9, align 8
 ; CHECK-NEXT:    %10 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 1
 ; CHECK-NEXT:    store i64 4, i64* %10, align 8
-; CHECK-NEXT:    call void @__tgt_target_data_begin(i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0))
+
+; CHECK-NEXT:    %handle = call %struct.__tgt_async_info @__tgt_target_data_begin_mapper_issue(i64 -1, i32 2, i8** %1, i8** %3, i64* %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.7, i64 0, i64 0), i8** null)
+; CHECK-NEXT:    call void @__tgt_target_data_begin_mapper_wait(i64 -1, %struct.__tgt_async_info %handle)
+
 ; CHECK-NEXT:    %11 = load i32, i32* %size.addr, align 4
 ; CHECK-NEXT:    %size.casted = zext i32 %11 to i64
 ; CHECK-NEXT:    %12 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 0
 ; CHECK-NEXT:    %13 = bitcast [2 x i8*]* %.offload_baseptrs2 to i64*
 ; CHECK-NEXT:    store i64 %size.casted, i64* %13, align 8
 ; CHECK-NEXT:    %14 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 0
 ; CHECK-NEXT:    %15 = bitcast [2 x i8*]* %.offload_ptrs3 to i64*
 ; CHECK-NEXT:    store i64 %size.casted, i64* %15, align 8
 ; CHECK-NEXT:    %16 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 1
 ; CHECK-NEXT:    %17 = bitcast i8** %16 to double**
 ; CHECK-NEXT:    store double* %a, double** %17, align 8
 ; CHECK-NEXT:    %18 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 1
 ; CHECK-NEXT:    %19 = bitcast i8** %18 to double**
 ; CHECK-NEXT:    store double* %a, double** %19, align 8
-; CHECK-NEXT:    %20 = call i32 @__tgt_target_teams(i64 -1, i8* nonnull @.__omp_offloading_heavyComputation2.region_id, i32 2, i8** nonnull %12, i8** nonnull %14, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.2, i64 0, i64 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.4, i64 0, i64 0), i32 0, i32 0)
-; CHECK-NEXT:    %21 = icmp eq i32 %20, 0
-; CHECK-NEXT:    br i1 %21, label %omp_offload.cont, label %omp_offload.failed
+; CHECK-NEXT:    %20 = call i32 @__tgt_target_teams_mapper(i64 -1, i8* nonnull @.__omp_offloading_heavyComputation2_region_id, i32 2, i8** nonnull %12, i8** nonnull %14, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.9, i64 0, i64 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.10, i64 0, i64 0), i8** null, i32 0, i32 0)
+; CHECK-NEXT:    %.not = icmp eq i32 %20, 0
+; CHECK-NEXT:    br i1 %.not, label %omp_offload.cont, label %omp_offload.failed
 ; CHECK:       omp_offload.failed:
 ; CHECK-NEXT:    call void @heavyComputation2FallBack(i64 %size.casted, double* %a)
 ; CHECK-NEXT:    br label %omp_offload.cont
 ; CHECK:       omp_offload.cont:
 ; CHECK-NEXT:    %rem = srem i32 %call, 7
-; CHECK-NEXT:    call void @__tgt_target_data_end(i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0))
+; CHECK-NEXT:    call void @__tgt_target_data_end_mapper(i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.7, i64 0, i64 0), i8** null)
 ; CHECK-NEXT:    ret i32 %rem
 ;
 entry:
   %size.addr = alloca i32, align 4
   %.offload_baseptrs = alloca [2 x i8*], align 8
   %.offload_ptrs = alloca [2 x i8*], align 8
   %.offload_sizes = alloca [2 x i64], align 8
   %.offload_baseptrs2 = alloca [2 x i8*], align 8
   %.offload_ptrs3 = alloca [2 x i8*], align 8
 
-  ; FIXME: Should have after splitting the runtime call __tgt_target_data_begin.
-  ; %device_id1 = alloca i64, align 8
-  ; %async_info1 = alloca %struct.__tgt_async_info, align 8
-
   store i32 %size, i32* %size.addr, align 4
-  %call = call i32 @rand()
+  %call = tail call i32 (...) @rand()
 
-  ; FIXME: This setup for the runtime call __tgt_target_data_begin should be
-  ;        split into its "issue" and "wait" counterpars. Here though, the "issue"
-  ;        cannot be moved upwards because it's not guaranteed that rand()
-  ;        won't modify *a. Nevertheless, the "wait" can be moved downwards.
-  ;        The call should be replaced to something like ...
-  ; Issue - this can't be moved upwards, *a might have aliases.
-  ; ... setup code ...
-  ; store i64 -1, i64* %device_id1, align 8
-  ; %handle1 = call i8* @__tgt_target_data_begin(i64* dereferenceable(8) %device_id1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0))
-  ; Wait - this is moved downards.
-  ; %device_id1_copy = load i64, i64* %device_id1, align 8 ; device_id
-  ; %queue1 = getelementptr inbounds %struct.__tgt_async_info, %struct.__tgt_async_info* %async_info1, i32 0, i32 0
-  ; store i8* %handle1, i8** %queue1, align 8
-  ; call void @__tgt_target_data_begin_wait(i64 %device_id1_copy, %struct.__tgt_async_info* dereferenceable(8) %async_info1)
   %conv = zext i32 %size to i64
   %0 = shl nuw nsw i64 %conv, 3
   %1 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 0
   %2 = bitcast [2 x i8*]* %.offload_baseptrs to double**
   store double* %a, double** %2, align 8
   %3 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 0
   %4 = bitcast [2 x i8*]* %.offload_ptrs to double**
   store double* %a, double** %4, align 8
   %5 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 0
   store i64 %0, i64* %5, align 8
   %6 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 1
   %7 = bitcast i8** %6 to i32**
   store i32* %size.addr, i32** %7, align 8
   %8 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 1
   %9 = bitcast i8** %8 to i32**
   store i32* %size.addr, i32** %9, align 8
   %10 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 1
   store i64 4, i64* %10, align 8
-  call void @__tgt_target_data_begin(i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0))
+  call void @__tgt_target_data_begin_mapper(i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.7, i64 0, i64 0), i8** null)
 
   %11 = load i32, i32* %size.addr, align 4
   %size.casted = zext i32 %11 to i64
   %12 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 0
   %13 = bitcast [2 x i8*]* %.offload_baseptrs2 to i64*
   store i64 %size.casted, i64* %13, align 8
   %14 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 0
   %15 = bitcast [2 x i8*]* %.offload_ptrs3 to i64*
   store i64 %size.casted, i64* %15, align 8
   %16 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 1
   %17 = bitcast i8** %16 to double**
   store double* %a, double** %17, align 8
   %18 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 1
   %19 = bitcast i8** %18 to double**
   store double* %a, double** %19, align 8
 
-  ; FIXME: The "wait" should be moved here.
-  %20 = call i32 @__tgt_target_teams(i64 -1, i8* nonnull @.__omp_offloading_heavyComputation2.region_id, i32 2, i8** nonnull %12, i8** nonnull %14, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.2, i64 0, i64 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.4, i64 0, i64 0), i32 0, i32 0)
-
-  %21 = icmp eq i32 %20, 0
-  br i1 %21, label %omp_offload.cont, label %omp_offload.failed
+  ; FIXME: call i8* @__tgt_target_data_begin_mapper_wait(...) should be moved here.
+  %20 = call i32 @__tgt_target_teams_mapper(i64 -1, i8* nonnull @.__omp_offloading_heavyComputation2_region_id, i32 2, i8** nonnull %12, i8** nonnull %14, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.9, i64 0, i64 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.10, i64 0, i64 0), i8** null, i32 0, i32 0)
+  %.not = icmp eq i32 %20, 0
+  br i1 %.not, label %omp_offload.cont, label %omp_offload.failed
 
 omp_offload.failed:                               ; preds = %entry
   call void @heavyComputation2FallBack(i64 %size.casted, double* %a)
   br label %omp_offload.cont
 
-omp_offload.cont:                                 ; preds = %entry, %omp_offload.failed
+omp_offload.cont:                                 ; preds = %omp_offload.failed, %entry
   %rem = srem i32 %call, 7
-  call void @__tgt_target_data_end(i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0))
+  call void @__tgt_target_data_end_mapper(i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.7, i64 0, i64 0), i8** null)
   ret i32 %rem
 }
 
 define internal void @heavyComputation2FallBack(i64 %size, double* %a) {
 entry:
   ; Fallback for offloading function heavyComputation2.
   ret void
 }

 ; CHECK-NEXT:  entry:
 ; CHECK-NEXT:    %size.addr = alloca i32, align 4
 ; CHECK-NEXT:    %.offload_baseptrs = alloca [2 x i8*], align 8
 ; CHECK-NEXT:    %.offload_ptrs = alloca [2 x i8*], align 8
 ; CHECK-NEXT:    %.offload_sizes = alloca [2 x i64], align 8
 ; CHECK-NEXT:    %.offload_baseptrs2 = alloca [2 x i8*], align 8
 ; CHECK-NEXT:    %.offload_ptrs3 = alloca [2 x i8*], align 8
 ; CHECK-NEXT:    store i32 %size, i32* %size.addr, align 4
-; CHECK-NEXT:    %call = call i32 @rand()
+; CHECK-NEXT:    %call = tail call i32 (...) @rand()
 ; CHECK-NEXT:    %conv = zext i32 %size to i64
 ; CHECK-NEXT:    %0 = shl nuw nsw i64 %conv, 3
 ; CHECK-NEXT:    %1 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 0
 ; CHECK-NEXT:    %2 = bitcast [2 x i8*]* %.offload_baseptrs to double**
 ; CHECK-NEXT:    store double* %a, double** %2, align 8
 ; CHECK-NEXT:    %3 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 0
 ; CHECK-NEXT:    %4 = bitcast [2 x i8*]* %.offload_ptrs to double**
 ; CHECK-NEXT:    store double* %a, double** %4, align 8
 ; CHECK-NEXT:    %5 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 0
 ; CHECK-NEXT:    store i64 %0, i64* %5, align 8
 ; CHECK-NEXT:    %6 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 1
 ; CHECK-NEXT:    %7 = bitcast i8** %6 to i32**
 ; CHECK-NEXT:    store i32* %size.addr, i32** %7, align 8
 ; CHECK-NEXT:    %8 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 1
 ; CHECK-NEXT:    %9 = bitcast i8** %8 to i32**
 ; CHECK-NEXT:    store i32* %size.addr, i32** %9, align 8
 ; CHECK-NEXT:    %10 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 1
 ; CHECK-NEXT:    store i64 4, i64* %10, align 8
-; CHECK-NEXT:    call void @__tgt_target_data_begin(i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0))
+
+; CHECK-NEXT:    %handle = call %struct.__tgt_async_info @__tgt_target_data_begin_mapper_issue(i64 -1, i32 2, i8** %1, i8** %3, i64* %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.7, i64 0, i64 0), i8** null)
+; CHECK-NEXT:    call void @__tgt_target_data_begin_mapper_wait(i64 -1, %struct.__tgt_async_info %handle)
+
 ; CHECK-NEXT:    %11 = load i32, i32* %size.addr, align 4
 ; CHECK-NEXT:    %size.casted = zext i32 %11 to i64
 ; CHECK-NEXT:    %12 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 0
 ; CHECK-NEXT:    %13 = bitcast [2 x i8*]* %.offload_baseptrs2 to i64*
 ; CHECK-NEXT:    store i64 %size.casted, i64* %13, align 8
 ; CHECK-NEXT:    %14 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 0
 ; CHECK-NEXT:    %15 = bitcast [2 x i8*]* %.offload_ptrs3 to i64*
 ; CHECK-NEXT:    store i64 %size.casted, i64* %15, align 8
 ; CHECK-NEXT:    %16 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 1
 ; CHECK-NEXT:    %17 = bitcast i8** %16 to double**
 ; CHECK-NEXT:    store double* %a, double** %17, align 8
 ; CHECK-NEXT:    %18 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 1
 ; CHECK-NEXT:    %19 = bitcast i8** %18 to double**
 ; CHECK-NEXT:    store double* %a, double** %19, align 8
-; CHECK-NEXT:    %20 = call i32 @__tgt_target_teams(i64 -1, i8* nonnull @.__omp_offloading_heavyComputation3.region_id, i32 2, i8** nonnull %12, i8** nonnull %14, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.2, i64 0, i64 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.4, i64 0, i64 0), i32 0, i32 0)
-; CHECK-NEXT:    %21 = icmp eq i32 %20, 0
-; CHECK-NEXT:    br i1 %21, label %omp_offload.cont, label %omp_offload.failed
+; CHECK-NEXT:    %20 = call i32 @__tgt_target_teams_mapper(i64 -1, i8* nonnull @.__omp_offloading_heavyComputation3_region_id, i32 2, i8** nonnull %12, i8** nonnull %14, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.9, i64 0, i64 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.10, i64 0, i64 0), i8** null, i32 0, i32 0)
+; CHECK-NEXT:    %.not = icmp eq i32 %20, 0
+; CHECK-NEXT:    br i1 %.not, label %omp_offload.cont, label %omp_offload.failed
 ; CHECK:       omp_offload.failed:
 ; CHECK-NEXT:    call void @heavyComputation3FallBack(i64 %size.casted, double* %a)
 ; CHECK-NEXT:    br label %omp_offload.cont
 ; CHECK:       omp_offload.cont:
 ; CHECK-NEXT:    %rem = srem i32 %call, 7
-; CHECK-NEXT:    call void @__tgt_target_data_end(i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0))
+; CHECK-NEXT:    call void @__tgt_target_data_end_mapper(i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.7, i64 0, i64 0), i8** null)
 ; CHECK-NEXT:    ret i32 %rem
 ;
 entry:
   %size.addr = alloca i32, align 4
   %.offload_baseptrs = alloca [2 x i8*], align 8
   %.offload_ptrs = alloca [2 x i8*], align 8
   %.offload_sizes = alloca [2 x i64], align 8
   %.offload_baseptrs2 = alloca [2 x i8*], align 8
   %.offload_ptrs3 = alloca [2 x i8*], align 8
-
-  ; FIXME: Should have after splitting the runtime call __tgt_target_data_begin.
-  ; %device_id1 = alloca i64, align 8
-  ; %async_info1 = alloca %struct.__tgt_async_info, align 8
-
   store i32 %size, i32* %size.addr, align 4
 
-  ; FIXME: The "issue" should be moved here.
-  %call = call i32 @rand()
+  ; FIXME: call i8* @__tgt_target_data_begin_mapper_issue(...) should be moved here.
+  %call = tail call i32 (...) @rand()
 
-  ; FIXME: This setup for the runtime call __tgt_target_data_begin should be
-  ;        split into its "issue" and "wait" counterpars and moved upwards
-  ;        and downwards, respectively. The call should be replaced to something
-  ;        like ...
-  ; Issue - this is moved upwards.
-  ; ... setup code ...
-  ; store i64 -1, i64* %device_id1, align 8
-  ; %handle1 = call i8* @__tgt_target_data_begin(i64* dereferenceable(8) %device_id1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0))
-  ; Wait - this is moved downards.
-  ; %device_id1_copy = load i64, i64* %device_id1, align 8 ; device_id
-  ; %queue1 = getelementptr inbounds %struct.__tgt_async_info, %struct.__tgt_async_info* %async_info1, i32 0, i32 0
-  ; store i8* %handle1, i8** %queue1, align 8
-  ; call void @__tgt_target_data_begin_wait(i64 %device_id1_copy, %struct.__tgt_async_info* dereferenceable(8) %async_info1)
   %conv = zext i32 %size to i64
   %0 = shl nuw nsw i64 %conv, 3
   %1 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 0
   %2 = bitcast [2 x i8*]* %.offload_baseptrs to double**
   store double* %a, double** %2, align 8
   %3 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 0
   %4 = bitcast [2 x i8*]* %.offload_ptrs to double**
   store double* %a, double** %4, align 8
   %5 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 0
   store i64 %0, i64* %5, align 8
   %6 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 1
   %7 = bitcast i8** %6 to i32**
   store i32* %size.addr, i32** %7, align 8
   %8 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 1
   %9 = bitcast i8** %8 to i32**
   store i32* %size.addr, i32** %9, align 8
   %10 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 1
   store i64 4, i64* %10, align 8
-  call void @__tgt_target_data_begin(i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0))
+  call void @__tgt_target_data_begin_mapper(i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.7, i64 0, i64 0), i8** null)
 
   %11 = load i32, i32* %size.addr, align 4
   %size.casted = zext i32 %11 to i64
   %12 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 0
   %13 = bitcast [2 x i8*]* %.offload_baseptrs2 to i64*
   store i64 %size.casted, i64* %13, align 8
   %14 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 0
   %15 = bitcast [2 x i8*]* %.offload_ptrs3 to i64*
   store i64 %size.casted, i64* %15, align 8
   %16 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 1
   %17 = bitcast i8** %16 to double**
   store double* %a, double** %17, align 8
   %18 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 1
   %19 = bitcast i8** %18 to double**
   store double* %a, double** %19, align 8
 
-  ; FIXME: The "wait" should be moved here.
-  %20 = call i32 @__tgt_target_teams(i64 -1, i8* nonnull @.__omp_offloading_heavyComputation3.region_id, i32 2, i8** nonnull %12, i8** nonnull %14, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.2, i64 0, i64 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.4, i64 0, i64 0), i32 0, i32 0)
-
-  %21 = icmp eq i32 %20, 0
-  br i1 %21, label %omp_offload.cont, label %omp_offload.failed
+  ; FIXME: call i8* @__tgt_target_data_begin_mapper_wait(...) should be moved here.
+  %20 = call i32 @__tgt_target_teams_mapper(i64 -1, i8* nonnull @.__omp_offloading_heavyComputation3_region_id, i32 2, i8** nonnull %12, i8** nonnull %14, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.9, i64 0, i64 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.10, i64 0, i64 0), i8** null, i32 0, i32 0)
+  %.not = icmp eq i32 %20, 0
+  br i1 %.not, label %omp_offload.cont, label %omp_offload.failed
 
 omp_offload.failed:                               ; preds = %entry
   call void @heavyComputation3FallBack(i64 %size.casted, double* %a)
   br label %omp_offload.cont
 
-omp_offload.cont:                                 ; preds = %entry, %omp_offload.failed
+omp_offload.cont:                                 ; preds = %omp_offload.failed, %entry
   %rem = srem i32 %call, 7
-  call void @__tgt_target_data_end(i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0))
+  call void @__tgt_target_data_end_mapper(i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.7, i64 0, i64 0), i8** null)
   ret i32 %rem
 }
 
 define internal void @heavyComputation3FallBack(i64 %size, double* %a) {
 entry:
   ; Fallback for offloading function heavyComputation3.
   ret void
 }

 ;  return random;
 ;}
 define dso_local i32 @dataTransferOnly1(double* noalias %a, i32 %size) {
 ; CHECK-LABEL: define {{[^@]+}}@dataTransferOnly1(double* noalias %a, i32 %size)
 ; CHECK-NEXT:  entry:
 ; CHECK-NEXT:    %.offload_baseptrs = alloca [1 x i8*], align 8
 ; CHECK-NEXT:    %.offload_ptrs = alloca [1 x i8*], align 8
 ; CHECK-NEXT:    %.offload_sizes = alloca [1 x i64], align 8
-; CHECK-NEXT:    %call = call i32 @rand()
+; CHECK-NEXT:    %call = tail call i32 (...) @rand()
 ; CHECK-NEXT:    %conv = zext i32 %size to i64
 ; CHECK-NEXT:    %0 = shl nuw nsw i64 %conv, 3
 ; CHECK-NEXT:    %1 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_baseptrs, i64 0, i64 0
 ; CHECK-NEXT:    %2 = bitcast [1 x i8*]* %.offload_baseptrs to double**
 ; CHECK-NEXT:    store double* %a, double** %2, align 8
 ; CHECK-NEXT:    %3 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i64 0, i64 0
 ; CHECK-NEXT:    %4 = bitcast [1 x i8*]* %.offload_ptrs to double**
 ; CHECK-NEXT:    store double* %a, double** %4, align 8
 ; CHECK-NEXT:    %5 = getelementptr inbounds [1 x i64], [1 x i64]* %.offload_sizes, i64 0, i64 0
 ; CHECK-NEXT:    store i64 %0, i64* %5, align 8
-; CHECK-NEXT:    call void @__tgt_target_data_begin(i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.5, i64 0, i64 0))
+
+; CHECK-NEXT:    %handle = call %struct.__tgt_async_info @__tgt_target_data_begin_mapper_issue(i64 -1, i32 1, i8** %1, i8** %3, i64* %5, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.11, i64 0, i64 0), i8** null)
+; CHECK-NEXT:    call void @__tgt_target_data_begin_mapper_wait(i64 -1, %struct.__tgt_async_info %handle)
+
 ; CHECK-NEXT:    %rem = urem i32 %call, %size
-; CHECK-NEXT:    call void @__tgt_target_data_end(i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.5, i64 0, i64 0))
+; CHECK-NEXT:    call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.11, i64 0, i64 0), i8** null)
 ; CHECK-NEXT:    ret i32 %rem
 ;
 entry:
   %.offload_baseptrs = alloca [1 x i8*], align 8
   %.offload_ptrs = alloca [1 x i8*], align 8
   %.offload_sizes = alloca [1 x i64], align 8
 
-  ; FIXME: Should have after splitting the runtime call __tgt_target_data_begin.
-  ; %device_id1 = alloca i64, align 8
-  ; %async_info1 = alloca %struct.__tgt_async_info, align 8
-
-  ; FIXME: The "issue" should be moved here.
-  %call = call i32 @rand()
-
-  ; FIXME: This setup for the runtime call __tgt_target_data_begin should be
-  ;        split into its "issue" and "wait" counterpars and moved upwards
-  ;        and downwards, respectively. The call should be replaced to something
-  ;        like ...
-  ; Issue - this is moved upwards.
-  ; ... setup code ...
-  ; store i64 -1, i64* %device_id1, align 8
-  ; %handle1 = call i8* @__tgt_target_data_begin(i64* dereferenceable(8) %device_id1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.5, i64 0, i64 0))
-  ; Wait - this is moved downards.
-  ; %device_id1_copy = load i64, i64* %device_id1, align 8 ; device_id
-  ; %queue1 = getelementptr inbounds %struct.__tgt_async_info, %struct.__tgt_async_info* %async_info1, i32 0, i32 0
-  ; store i8* %handle1, i8** %queue1, align 8
-  ; call void @__tgt_target_data_begin_wait(i64 %device_id1_copy, %struct.__tgt_async_info* dereferenceable(8) %async_info1)
+  ; FIXME: call i8* @__tgt_target_data_begin_issue_mapper(...) should be moved here.
+  %call = tail call i32 (...) @rand()
+
   %conv = zext i32 %size to i64
   %0 = shl nuw nsw i64 %conv, 3
   %1 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_baseptrs, i64 0, i64 0
   %2 = bitcast [1 x i8*]* %.offload_baseptrs to double**
   store double* %a, double** %2, align 8
   %3 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i64 0, i64 0
   %4 = bitcast [1 x i8*]* %.offload_ptrs to double**
   store double* %a, double** %4, align 8
   %5 = getelementptr inbounds [1 x i64], [1 x i64]* %.offload_sizes, i64 0, i64 0
   store i64 %0, i64* %5, align 8
-  call void @__tgt_target_data_begin(i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.5, i64 0, i64 0))
+  call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.11, i64 0, i64 0), i8** null)
 
   %rem = urem i32 %call, %size
 
-  ; FIXME: The "wait" should be moved here.
-  call void @__tgt_target_data_end(i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.5, i64 0, i64 0))
+  call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.11, i64 0, i64 0), i8** null)
   ret i32 %rem
 }
 
-declare dso_local void @__tgt_target_data_begin(i64, i32, i8**, i8**, i64*, i64*)
-
-declare dso_local i32 @__tgt_target_teams(i64, i8*, i32, i8**, i8**, i64*, i64*, i32, i32)
-
-declare dso_local void @__tgt_target_data_end(i64, i32, i8**, i8**, i64*, i64*)
+; CHECK: declare %struct.__tgt_async_info @__tgt_target_data_begin_mapper_issue(i64, i32, i8**, i8**, i64*, i64*, i8**)
+; CHECK: declare void @__tgt_target_data_begin_mapper_wait(i64, %struct.__tgt_async_info)
 
-declare dso_local i32 @rand()
+declare void @__tgt_target_data_begin_mapper(i64, i32, i8**, i8**, i64*, i64*, i8**)
+declare i32 @__tgt_target_teams_mapper(i64, i8*, i32, i8**, i8**, i64*, i64*, i8**, i32, i32)
+declare void @__tgt_target_data_end_mapper(i64, i32, i8**, i8**, i64*, i64*, i8**)
 
-; FIXME: These two function declarations must be generated after splitting the runtime function
-;        __tgt_target_data_begin.
-; declare dso_local i8* @__tgt_target_data_begin_issue(i64* dereferenceable(8), i32, i8**, i8**, i64*, i64*)
-; declare dso_local void @__tgt_target_data_begin_wait(i64, %struct.__tgt_async_info* dereferenceable(8))
+declare dso_local i32 @rand(...)