[OpenMP] Fix `present` for exit from `omp target data`

[jdenny]

Summary

Without this patch, the following example fails but shouldn't
according to OpenMP TR8:

#pragma omp target enter data map(alloc:i)
#pragma omp target data map(present, alloc: i)
{
  #pragma omp target exit data map(delete:i)
} // fails presence check here

OpenMP TR8 sec. 2.22.7.1 "map Clause", p. 321, L23-26 states:

If the map clause appears on a target, target data, target enter
data or target exit data construct with a present map-type-modifier
then on entry to the region if the corresponding list item does not
appear in the device data environment an error occurs and the
program terminates.

There is no corresponding statement about the exit from a region.
Thus, the present modifier should:

1. Check for presence upon entry into any region, including a target exit data region. This behavior is already implemented correctly.

2. Should not check for presence upon exit from any region, including a target or target data region. Without this patch, this behavior is not implemented correctly, breaking the above example.

In the case of target data, this patch fixes the latter behavior by
removing the present modifier from the map types Clang generates for
the runtime call at the end of the region.

In the case of target, we have not found a valid OpenMP program for
which such a fix would matter. It appears that, if a program can
guarantee that data is present at the beginning of a target region
so that there's no error there, that data is also guaranteed to be
present at the end.

[grokos]

What confuses me about this interpretation of the standard is the inconsistency at data exit. So if we have an explicit omp target exit data map(present...) then we should respect the "present" semantics, whereas when we have a scoped data exit:

#pragma omp target data map(present,...)
{
  ...
} // implicit "exit data" here

then "present" should be ignored.

I agree that the paragraph from the standard leaves little room for other interpretations, I'd just like to point out that it looks inconsistent - at least to me.

[jdenny]

I don't know if the OpenMP committee has any documented rationale for this behavior. I can say that the OpenACC committee is considering the same semantics. However, the issues to consider are not identical. For example, OpenACC has a separate structured reference counter, meaning it should be impossible for such data not to be present at the exit of a data construct unless you've shut down the runtime.

[RaviNarayanaswamy]

In D84422#2170285, @grokos wrote:

What confuses me about this interpretation of the standard is the inconsistency at data exit. So if we have an explicit omp target exit data map(present...) then we should respect the "present" semantics, whereas when we have a scoped data exit:

#pragma omp target data map(present,...)
{
  ...
} // implicit "exit data" here

then "present" should be ignored.

I agree that the paragraph from the standard leaves little room for other interpretations, I'd just like to point out that it looks inconsistent - at least to me.

When you use present on a variable on a scoped target data region, you cannot delete that object in the scope. I would say this is a test case error. It should still be present on exit, checking for that is maybe redundant

[grokos]

So is the test case that motivated this patch illegal OpenMP code?

#pragma omp target enter data map(alloc:i)
#pragma omp target data map(present, alloc: i)
{
  #pragma omp target exit data map(delete:i) // you cannot delete that object in the scope, illegal code?
} // fails presence check here

[RaviNarayanaswamy]

In D84422#2170667, @grokos wrote:

So is the test case that motivated this patch illegal OpenMP code?

#pragma omp target enter data map(alloc:i)
#pragma omp target data map(present, alloc: i)
{
  #pragma omp target exit data map(delete:i) // you cannot delete that object in the scope, illegal code?
} // fails presence check here

According to spec the test should work. ie should not check for presence on exit from a blocked openmp pragma scope.

[jdenny]

In D84422#2170702, @RaviNarayanaswamy wrote:

In D84422#2170667, @grokos wrote:

So is the test case that motivated this patch illegal OpenMP code?

#pragma omp target enter data map(alloc:i)
#pragma omp target data map(present, alloc: i)
{
  #pragma omp target exit data map(delete:i) // you cannot delete that object in the scope, illegal code?
} // fails presence check here

According to spec the test should work. ie should not check for presence on exit from a blocked openmp pragma scope.

It sounds like this patch's motivation is correct then. Has anyone clarified the motivation for this behavior?

[grokos]

So let's proceed with the patch.

Instead of introducing new API functions and making all these changes in all these files, wouldn't it be easier if we just unset the PRESENT flag from arg_types in clang when we generate the call to __tgt_target_data_end_* if we are exiting from a scoped environment?

[jdenny]

In D84422#2172898, @jdenny wrote:

Has anyone clarified the motivation for this behavior?

I meant, is there any insight into why the spec specifies this behavior?

In D84422#2172926, @grokos wrote:

Instead of introducing new API functions and making all these changes in all these files, wouldn't it be easier if we just unset the PRESENT flag from arg_types in clang when we generate the call to __tgt_target_data_end_* if we are exiting from a scoped environment?

Ah, that does sound simpler. Thanks. I'll look into it.

Suppressing the presence check on exit from omp target would require a runtime change in addition to the Clang change you suggest for omp target data. However, I've so far failed to formulate a reasonable test case. Specifically, I don't yet see a way to guarantee that the data will definitely be present at the start of omp target but might not be present by the end. Is it possible? If not, then maybe we should leave the check in place for omp target.

[RaviNarayanaswamy]

In D84422#2173372, @jdenny wrote:

In D84422#2172898, @jdenny wrote:

Has anyone clarified the motivation for this behavior?

I meant, is there any insight into why the spec specifies this behavior?

In D84422#2172926, @grokos wrote:

Instead of introducing new API functions and making all these changes in all these files, wouldn't it be easier if we just unset the PRESENT flag from arg_types in clang when we generate the call to __tgt_target_data_end_* if we are exiting from a scoped environment?

Ah, that does sound simpler. Thanks. I'll look into it.

Suppressing the presence check on exit from omp target would require a runtime change in addition to the Clang change you suggest for omp target data. However, I've so far failed to formulate a reasonable test case. Specifically, I don't yet see a way to guarantee that the data will definitely be present at the start of omp target but might not be present by the end. Is it possible? If not, then maybe we should leave the check in place for omp target.

I would rather not have a check if not required by the spec as it would an unnecessary overhead to performance.

[jdenny]

In D84422#2173449, @RaviNarayanaswamy wrote:

In D84422#2173372, @jdenny wrote:

In D84422#2172898, @jdenny wrote:

Has anyone clarified the motivation for this behavior?

I meant, is there any insight into why the spec specifies this behavior?

In D84422#2172926, @grokos wrote:

Instead of introducing new API functions and making all these changes in all these files, wouldn't it be easier if we just unset the PRESENT flag from arg_types in clang when we generate the call to __tgt_target_data_end_* if we are exiting from a scoped environment?

Ah, that does sound simpler. Thanks. I'll look into it.

Suppressing the presence check on exit from omp target would require a runtime change in addition to the Clang change you suggest for omp target data. However, I've so far failed to formulate a reasonable test case. Specifically, I don't yet see a way to guarantee that the data will definitely be present at the start of omp target but might not be present by the end. Is it possible? If not, then maybe we should leave the check in place for omp target.

I would rather not have a check if not required by the spec as it would an unnecessary overhead to performance.

I've added a comment to the runtime code that performs the check. As you can see, the check is performed regardless. It's just a question of whether the runtime treats it as an error. I don't think performance is an issue.

My concern here is that it will be hard to justify changes to the runtime if I cannot formulate a use case.

openmp/libomptarget/src/omptarget.cpp

511

This is where the runtime performs the check.

[grokos]

In D84422#2173500, @jdenny wrote:

I've added a comment to the runtime code that performs the check. As you can see, the check is performed regardless. It's just a question of whether the runtime treats it as an error. I don't think performance is an issue.

My concern here is that it will be hard to justify changes to the runtime if I cannot formulate a use case.

Thinking about it, I don't think there can be a case where something is present upon entering a target region and not be present when we're exiting. Whatever code comprises the target region is code executed on the device - it cannot modify the state of host objects (i.e. libomptarget) in any possible way. E.g. the kernel cannot invoke libomptarget functions, allocate memory, map/unmap data etc.

The only case where something like this would be possible is if we have multiple host threads executing async offloading. In such a case, one thread may launch a target region at a moment when the requested mapping is present on the device and while the kernel is executing some other thread performs a target data exit on the desired mapping. Upon exiting the kernel, the mapping will no longer be present but this is clearly a race condition (user's fault), so I don't think we should pay attention to such a scenario.

[jdenny]

In D84422#2176802, @grokos wrote:

In D84422#2173500, @jdenny wrote:

I've added a comment to the runtime code that performs the check. As you can see, the check is performed regardless. It's just a question of whether the runtime treats it as an error. I don't think performance is an issue.

My concern here is that it will be hard to justify changes to the runtime if I cannot formulate a use case.

Thinking about it, I don't think there can be a case where something is present upon entering a target region and not be present when we're exiting. Whatever code comprises the target region is code executed on the device - it cannot modify the state of host objects (i.e. libomptarget) in any possible way. E.g. the kernel cannot invoke libomptarget functions, allocate memory, map/unmap data etc.

The only case where something like this would be possible is if we have multiple host threads executing async offloading. In such a case, one thread may launch a target region at a moment when the requested mapping is present on the device and while the kernel is executing some other thread performs a target data exit on the desired mapping. Upon exiting the kernel, the mapping will no longer be present but this is clearly a race condition (user's fault), so I don't think we should pay attention to such a scenario.

Exactly. As far as I can tell, the runtime simply needs a comment that explains this issue in the case of omp target.

[jdenny]

Rewrite patch as discussed: instead of generating different runtime calls for the end of an omp target data vs. the beginning of an omp target exit data so that the runtime can determine when to ignore present, change Clang to filter present from the map types at the end of an omp target data.

[grokos]

This looks much better now. I don't have any other comments. Since this patch is now essentially a clang-only patch, I'll let @ABataev accept it or post comments.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8686

Can this new flag be encapsulated in Info?

[jdenny]

Replaced SeparateBeginEnd parameter with new TargetDataInfo field as requested. Rebased.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8843

Do not append param here, use the one from Info

[jdenny]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8843

Info.SeparateBeginEndCalls and ForEndCall do not represent the same thing. If Info.SeparateBeginEndCalls=true, as in emitTargetDataCalls below, then emitOffloadingArraysArgument is called twice with the same Info, once with ForEndCall=false and once with ForEndCall=true.

[ABataev]

LG

[jdenny]

Thanks for the review.

As discussed during the 7/29 call, I'll wait to push until we're sure about what the OpenMP committee intended here. I'm pursuing this and will report back when I have more information.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

 
 /// 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::MapCombinedInfoTy &CombinedInfo,
                      CGOpenMPRuntime::TargetDataInfo &Info) {
   CodeGenModule &CGM = CGF.CGM;
   ASTContext &Ctx = CGF.getContext();
 
   // Reset the array information.
   Info.clearArrayInfo();
   Info.NumberOfPtrs = CombinedInfo.BasePointers.size();
 
   if (Info.NumberOfPtrs) {
     // Detect if we have any capture size requiring runtime evaluation of the

         /*Idx0=*/0,
         /*Idx1=*/0);
     MappersArrayArg =
         Info.HasMapper
             ? CGF.Builder.CreatePointerCast(Info.MappersArray, CGM.VoidPtrPtrTy)
             : llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
   } 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.

                                      InputInfo.MappersArray.getPointer()};
 
     // Select the right runtime function call for each standalone
     // directive.
     const bool HasNowait = D.hasClausesOfKind<OMPNowaitClause>();
     RuntimeFunction RTLFn;
     switch (D.getDirectiveKind()) {
     case OMPD_target_enter_data:
-      RTLFn = HasNowait ? OMPRTL___tgt_target_data_begin_nowait_mapper
-                        : OMPRTL___tgt_target_data_begin_mapper;
+      RTLFn = HasNowait ? OMPRTL___tgt_target_enter_data_nowait_mapper
+                        : OMPRTL___tgt_target_enter_data_mapper;
       break;
     case OMPD_target_exit_data:
-      RTLFn = HasNowait ? OMPRTL___tgt_target_data_end_nowait_mapper
-                        : OMPRTL___tgt_target_data_end_mapper;
+      RTLFn = HasNowait ? OMPRTL___tgt_target_exit_data_nowait_mapper
+                        : OMPRTL___tgt_target_exit_data_mapper;
       break;
     case OMPD_target_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:

clang/test/OpenMP/declare_mapper_codegen.cpp

void foo(int a){
   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
   // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
   // CK0: call void [[KERNEL:@.+]](%class.C* [[VAL]])
   #pragma omp target teams map(mapper(id),to: c) nowait
   {
     ++c.a;
   }
 
-  // CK0-DAG: call void @__tgt_target_data_begin_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EDSIZES]]{{.+}}, {{.+}}[[EDTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
+  // CK0-DAG: call void @__tgt_target_enter_data_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EDSIZES]]{{.+}}, {{.+}}[[EDTYPES]]{{.+}}, 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]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
   // 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]], i[[sz]] 0, i[[sz]] 0
   // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
   // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
   // CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
   // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
   #pragma omp target enter data map(mapper(id),to: c)
 
-  // CK0-DAG: call void @__tgt_target_data_begin_nowait_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EDNWSIZES]]{{.+}}, {{.+}}[[EDNWTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
+  // CK0-DAG: call void @__tgt_target_enter_data_nowait_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EDNWSIZES]]{{.+}}, {{.+}}[[EDNWTYPES]]{{.+}}, 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]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
   // 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]], i[[sz]] 0, i[[sz]] 0
   // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
   // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
   // CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
   // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
   #pragma omp target enter data map(mapper(id),to: c) nowait
 
-  // CK0-DAG: call void @__tgt_target_data_end_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EXDSIZES]]{{.+}}, {{.+}}[[EXDTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
+  // CK0-DAG: call void @__tgt_target_exit_data_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EXDSIZES]]{{.+}}, {{.+}}[[EXDTYPES]]{{.+}}, 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]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
   // 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]], i[[sz]] 0, i[[sz]] 0
   // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
   // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
   // CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
   // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
   // CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
   #pragma omp target exit data map(mapper(id),from: c)
 
-  // CK0-DAG: call void @__tgt_target_data_end_nowait_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EXDNWSIZES]]{{.+}}, {{.+}}[[EXDNWTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
+  // CK0-DAG: call void @__tgt_target_exit_data_nowait_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EXDNWSIZES]]{{.+}}, {{.+}}[[EXDNWTYPES]]{{.+}}, 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]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
   // 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]], i[[sz]] 0, i[[sz]] 0
   // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
   // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**

clang/test/OpenMP/target_enter_data_codegen.cpp

 // CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1061]
 
 // CK1-LABEL: _Z3fooi
 void foo(int arg) {
   int la;
   float lb[arg];
 
   // Region 00
-  // CK1-DAG: call void @__tgt_target_data_begin_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+  // CK1-DAG: call void @__tgt_target_enter_data_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
   // CK1-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [100 x double]**
   // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [100 x double]**
   // CK1-DAG: store [100 x double]* @gc, [100 x double]** [[CBP0]]
   // CK1-DAG: store [100 x double]* @gc, [100 x double]** [[CP0]]
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
-  // CK1-NOT: __tgt_target_data_end
+  // CK1-NOT: __tgt_target_exit_data
   #pragma omp target enter data if(1+3-5) device(arg) map(alloc: gc) nowait
   {++arg;}
 
   // Region 01
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   #pragma omp target enter data map(to: la) if(1+3-4)
   {++arg;}
 
   // Region 02
   // CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
   // CK1: [[IFTHEN]]
-  // CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null)
+  // CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null)
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
   // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
   // CK1-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]]
   // CK1-DAG: store i32* [[VAR0]], i32** [[CP0]]
   // CK1: br label %[[IFEND:[^,]+]]
 
   // CK1: [[IFELSE]]
   // CK1: br label %[[IFEND]]
   // CK1: [[IFEND]]
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
-  // CK1-NOT: __tgt_target_data_end
+  // CK1-NOT: __tgt_target_exit_data
   #pragma omp target enter data map(to: arg) if(arg) device(4)
   {++arg;}
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   {++arg;}
 
   // Region 03
-  // CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** null)
+  // CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** null)
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
   // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
   // CK1-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
   // CK1-DAG: store float* [[VAR0]], float** [[CP0]]
   // CK1-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]]
   // CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
   // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
   // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
-  // CK1-NOT: __tgt_target_data_end
+  // CK1-NOT: __tgt_target_exit_data
   #pragma omp target enter data map(always, to: lb)
   {++arg;}
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   {++arg;}
 
   // Region 04
-  // CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** null)
+  // CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** null)
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
   // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double***
   // CK1-DAG: store [[ST]]* @gb, [[ST]]** [[CBP0]]
   // CK1-DAG: store double** getelementptr inbounds ([[ST]], [[ST]]* @gb, i32 0, i32 1), double*** [[CP0]]
 
 
   // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
   // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
   // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double***
   // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double**
   // CK1-DAG: store double** getelementptr inbounds ([[ST]], [[ST]]* @gb, i32 0, i32 1), double*** [[CBP1]]
   // CK1-DAG: store double* [[SEC1:%.+]], double** [[CP1]]
   // CK1-DAG: [[SEC1]] = getelementptr inbounds {{.+}}double* [[SEC11:%[^,]+]], i{{.+}} 0
   // CK1-DAG: [[SEC11]] = load double*, double** getelementptr inbounds ([[ST]], [[ST]]* @gb, i32 0, i32 1),
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
-  // CK1-NOT: __tgt_target_data_end
+  // CK1-NOT: __tgt_target_exit_data
   #pragma omp target enter data map(to: gb.b[:3])
   {++arg;}
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   {++arg;}
 
   // Region 05
-  // CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}, i8** null)
+  // CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}, i8** null)
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
   // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
   // CK1-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
   // CK1-DAG: store float* [[VAR0]], float** [[CP0]]
   // CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
   // CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
   // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
   // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
-  // CK1-NOT: __tgt_target_data_end
+  // CK1-NOT: __tgt_target_exit_data
   #pragma omp target enter data map(close, to: lb)
   {++arg;}
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   {++arg;}
 
   // Region 06
-  // CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}, i8** null)
+  // CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}, i8** null)
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
   // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
   // CK1-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
   // CK1-DAG: store float* [[VAR0]], float** [[CP0]]
   // CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
   // CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
   // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
   // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
-  // CK1-NOT: __tgt_target_data_end
+  // CK1-NOT: __tgt_target_exit_data
   #pragma omp target enter data map(always close, to: lb)
   {++arg;}
 }
 #endif
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK1A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1A --check-prefix CK1A-64
 // RUN: %clang_cc1 -DCK1A -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1A --check-prefix CK1A-64

 // CK1A: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1425]]]
 
 // CK1A-LABEL: _Z3fooi
 void foo(int arg) {
   int la;
   float lb[arg];
 
   // Region 00
-  // CK1A-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:32|64]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK1A-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:32|64]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
   // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
   // CK1A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
   // CK1A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
   // CK1A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
   // CK1A-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
   // CK1A-DAG: store float* [[VAR0]], float** [[CP0]]
   // CK1A-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
   // CK1A-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
   // CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
   // CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
   // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
-  // CK1A-NOT: __tgt_target_data_end
+  // CK1A-NOT: __tgt_target_exit_data
   #pragma omp target enter data map(present, to: lb)
   {++arg;}
 
   // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   {++arg;}
 
   // Region 01
-  // CK1A-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+  // CK1A-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
   // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
   // CK1A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
   // CK1A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
   // CK1A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
   // CK1A-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
   // CK1A-DAG: store float* [[VAR0]], float** [[CP0]]
   // CK1A-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
   // CK1A-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
   // CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
   // CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
   // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
-  // CK1A-NOT: __tgt_target_data_end
+  // CK1A-NOT: __tgt_target_exit_data
   #pragma omp target enter data map(always close present, to: lb)
   {++arg;}
 }
 #endif
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-64

 int bar(int arg){
   ST<int> A;
   return A.foo(arg);
 }
 
 // Region 00
 // CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
 // CK2: [[IFTHEN]]
-// CK2-DAG: call void @__tgt_target_data_begin_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+// CK2-DAG: call void @__tgt_target_enter_data_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
 // CK2-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
 // CK2-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
 // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 // CK2-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0

 // CK2-DAG: [[SEC111]] = getelementptr inbounds {{.*}}[[ST]]* [[VAR0]], i32 0, i32 1
 
 // CK2: br label %[[IFEND:[^,]+]]
 
 // CK2: [[IFELSE]]
 // CK2: br label %[[IFEND]]
 // CK2: [[IFEND]]
 // CK2: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
-// CK2-NOT: __tgt_target_data_end
+// CK2-NOT: __tgt_target_exit_data
 #endif
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64
 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-64
 // RUN: %clang_cc1 -DCK3 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-32
 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-32
 
 // RUN: %clang_cc1 -DCK3 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 // RUN: %clang_cc1 -DCK3 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 // SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
 #ifdef CK3
 
 // CK3-LABEL: no_target_devices
 void no_target_devices(int arg) {
-  // CK3-NOT: tgt_target_data_begin
+  // CK3-NOT: tgt_target_enter_data
   // CK3: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
-  // CK3-NOT: tgt_target_data_end
+  // CK3-NOT: tgt_target_exit_data
   // CK3: ret
   #pragma omp target enter data map(to: arg) if(arg) device(4)
   {++arg;}
 }
 #endif
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-64
 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s

 // RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck --check-prefix SIMD-ONLY4 %s
 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY4 %s
 // SIMD-ONLY4-NOT: {{__kmpc|__tgt}}
 #ifdef CK4
 
 // CK4-LABEL: device_side_scan
 void device_side_scan(int arg) {
-  // CK4: tgt_target_data_begin
+  // CK4: tgt_target_enter_data
   // CK4: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   // CK4: ret
-  // TCK4-NOT: tgt_target_data_begin
+  // TCK4-NOT: tgt_target_enter_data
   #pragma omp target enter data map(to: arg) if(arg) device(4)
   {++arg;}
 }
 #endif
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-64
 // RUN: %clang_cc1 -DCK5 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK5 --check-prefix CK5-64

 int bar(int arg){
   STT<int> A;
   return A.foo(arg);
 }
 
 // Region 00
 // CK5: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
 // CK5: [[IFTHEN]]
-// CK5-DAG: call void @__tgt_target_data_begin_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+// CK5-DAG: call void @__tgt_target_enter_data_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
 // CK5-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
 // CK5-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
 // CK5-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK5-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 // CK5-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
 // CK5-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 // CK5-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0

 // CK5-DAG: [[SEC111]] = getelementptr inbounds {{.*}}[[STT]]* [[VAR0]], i32 0, i32 1
 
 // CK5: br label %[[IFEND:[^,]+]]
 
 // CK5: [[IFELSE]]
 // CK5: br label %[[IFEND]]
 // CK5: [[IFEND]]
 // CK5: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
-// CK5-NOT: __tgt_target_data_end
+// CK5-NOT: __tgt_target_exit_data
 #endif
 #endif

clang/test/OpenMP/target_enter_data_depend_codegen.cpp

void foo(int arg) {
   // CK1: call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   // CK1: = call i32 [[TASK_ENTRY4]](i32 %{{.+}}, %struct.kmp_task_t_with_privates{{.+}}* [[RES_BC]])
   // CK1: call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   #pragma omp target enter data map(to:gb.b[:3]) depend(in: gb.b[:3], la, lb, gc, arg)
   {++arg;}
 }
 
 // CK1: define internal{{.*}} i32 [[TASK_ENTRY0]](i32{{.*}}, %struct.kmp_task_t_with_privates* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_begin_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_enter_data_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** [[GEPM:%.+]])
 // CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
 // CK1-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 // CK1-DAG: [[GEPM]] = getelementptr inbounds {{.+}}[[M:%[^,]+]]
 // CK1-DAG: [[BP]] = load [1 x i8*]*, [1 x i8*]** [[BP_PRIV:%[^,]+]],
 // CK1-DAG: [[P]] = load [1 x i8*]*, [1 x i8*]** [[P_PRIV:%[^,]+]],
 // CK1-DAG: [[S]] = load [1 x i64]*, [1 x i64]** [[S_PRIV:%[^,]+]],
 // CK1-DAG: [[M]] = load [1 x i8*]*, [1 x i8*]** [[M_PRIV:%[^,]+]],
 // CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [1 x i8*]** [[BP_PRIV]], [1 x i8*]** [[P_PRIV]], [1 x i64]** [[S_PRIV]], [1 x i8*]** [[M_PRIV]])
 // CK1: ret i32 0
 // CK1: }
 
 // CK1: define internal{{.*}} i32 [[TASK_ENTRY2]](i32{{.*}}, %struct.kmp_task_t_with_privates{{.+}}* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** [[GEPM:%.+]])
 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 // CK1-DAG: [[GEPM]] = getelementptr inbounds {{.+}}[[M:%[^,]+]]
 // CK1-DAG: [[BP]] = load [1 x i8*]*, [1 x i8*]** [[BP_PRIV:%[^,]+]],
 // CK1-DAG: [[P]] = load [1 x i8*]*, [1 x i8*]** [[P_PRIV:%[^,]+]],
 // CK1-DAG: [[S]] = load [1 x i64]*, [1 x i64]** [[S_PRIV:%[^,]+]],
 // CK1-DAG: [[M]] = load [1 x i8*]*, [1 x i8*]** [[M_PRIV:%[^,]+]],
 // CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [1 x i8*]** [[BP_PRIV]], [1 x i8*]** [[P_PRIV]], [1 x i64]** [[S_PRIV]], [1 x i8*]** [[M_PRIV]])
 // CK1: ret i32 0
 // CK1: }
 
 // CK1: define internal{{.*}} i32 [[TASK_ENTRY3]](i32{{.*}}, %struct.kmp_task_t_with_privates{{.+}}* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** [[GEPM:%.+]])
 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 // CK1-DAG: [[GEPM]] = getelementptr inbounds {{.+}}[[M:%[^,]+]]
 // CK1-DAG: [[BP]] = load [1 x i8*]*, [1 x i8*]** [[BP_PRIV:%[^,]+]],
 // CK1-DAG: [[P]] = load [1 x i8*]*, [1 x i8*]** [[P_PRIV:%[^,]+]],
 // CK1-DAG: [[S]] = load [1 x i64]*, [1 x i64]** [[S_PRIV:%[^,]+]],
 // CK1-DAG: [[M]] = load [1 x i8*]*, [1 x i8*]** [[M_PRIV:%[^,]+]],
 // CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [1 x i8*]** [[BP_PRIV]], [1 x i8*]** [[P_PRIV]], [1 x i64]** [[S_PRIV]], [1 x i8*]** [[M_PRIV]])
-// CK1-NOT: __tgt_target_data_end
+// CK1-NOT: __tgt_target_exit_data
 // CK1: ret i32 0
 // CK1: }
 
 // CK1: define internal{{.*}} i32 [[TASK_ENTRY4]](i32{{.*}}, %struct.kmp_task_t_with_privates{{.+}}* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** [[GEPM:%.+]])
 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 // CK1-DAG: [[GEPM]] = getelementptr inbounds {{.+}}[[M:%[^,]+]]
 // CK1-DAG: [[BP]] = load [2 x i8*]*, [2 x i8*]** [[BP_PRIV:%[^,]+]],
 // CK1-DAG: [[P]] = load [2 x i8*]*, [2 x i8*]** [[P_PRIV:%[^,]+]],
 // CK1-DAG: [[S]] = load [2 x i64]*, [2 x i64]** [[S_PRIV:%[^,]+]],
 // CK1-DAG: [[M]] = load [2 x i8*]*, [2 x i8*]** [[M_PRIV:%[^,]+]],
 // CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [2 x i8*]** [[BP_PRIV]], [2 x i8*]** [[P_PRIV]], [2 x i64]** [[S_PRIV]], [2 x i8*]** [[M_PRIV]])
-// CK1-NOT: __tgt_target_data_end
+// CK1-NOT: __tgt_target_exit_data
 // CK1: ret i32 0
 // CK1: }
 
 #endif

clang/test/OpenMP/target_exit_data_codegen.cpp

 // CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1062]
 
 // CK1-LABEL: _Z3fooi
 void foo(int arg) {
   int la;
   float lb[arg];
 
   // Region 00
-  // CK1-NOT: __tgt_target_data_begin
-  // CK1-DAG: call void @__tgt_target_data_end_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+  // CK1-NOT: __tgt_target_enter_data
+  // CK1-DAG: call void @__tgt_target_exit_data_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
   // CK1-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [100 x double]**
   // CK1-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to [100 x double]**
   // CK1-DAG: store [100 x double]* @gc, [100 x double]** [[BPC0]]
   // CK1-DAG: store [100 x double]* @gc, [100 x double]** [[PC0]]
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   #pragma omp target exit data if(1+3-5) device(arg) map(from: gc) nowait
   {++arg;}
 
   // Region 01
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   #pragma omp target exit data map(release: la) if(1+3-4)
   {++arg;}
 
   // Region 02
-  // CK1-NOT: __tgt_target_data_begin
+  // CK1-NOT: __tgt_target_enter_data
   // CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
   // CK1: [[IFTHEN]]
-  // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null)
+  // CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null)
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
   // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
   // CK1-DAG: store i32* [[VAL0:%[^,]+]], i32** [[CBP0]]
   // CK1-DAG: store i32* [[VAL0]], i32** [[CP0]]
   // CK1: br label %[[IFEND:[^,]+]]
 
   // CK1: [[IFELSE]]
   // CK1: br label %[[IFEND]]
   // CK1: [[IFEND]]
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   #pragma omp target exit data map(release: arg) if(arg) device(4)
   {++arg;}
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   {++arg;}
 
   // Region 03
-  // CK1-NOT: __tgt_target_data_begin
-  // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** null)
+  // CK1-NOT: __tgt_target_enter_data
+  // CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** null)
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
   // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
   // CK1-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]]
   // CK1-DAG: store float* [[VAL0]], float** [[CP0]]
   // CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
   // CK1-64-DAG: [[CSVAL0]] = mul nuw i[[sz]] %{{[^,]+}}, 4
   // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
   // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   #pragma omp target exit data map(always, from: lb)
   {++arg;}
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   {++arg;}
 
   // Region 04
-  // CK1-NOT: __tgt_target_data_begin
-  // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** null)
+  // CK1-NOT: __tgt_target_enter_data
+  // CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** null)
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
   // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double***
   // CK1-DAG: store [[ST]]* @gb, [[ST]]** [[CBP0]]
void foo(int arg) {
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   #pragma omp target exit data map(release: gb.b[:3])
   {++arg;}
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   {++arg;}
 
   // Region 05
-  // CK1-NOT: __tgt_target_data_begin
-  // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}, i8** null)
+  // CK1-NOT: __tgt_target_enter_data
+  // CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}, i8** null)
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
   // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
   // CK1-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]]
   // CK1-DAG: store float* [[VAL0]], float** [[CP0]]
   // CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
   // CK1-64-DAG: [[CSVAL0]] = mul nuw i[[sz]] %{{[^,]+}}, 4
   // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
   // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   #pragma omp target exit data map(close, from: lb)
   {++arg;}
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   {++arg;}
 
   // Region 06
-  // CK1-NOT: __tgt_target_data_begin
-  // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}, i8** null)
+  // CK1-NOT: __tgt_target_enter_data
+  // CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}, i8** null)
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**

 
 // CK2-LABEL: _Z3bari
 int bar(int arg){
   ST<int> A;
   return A.foo(arg);
 }
 
 // Region 00
-// CK2-NOT: __tgt_target_data_begin
+// CK2-NOT: __tgt_target_enter_data
 // CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
 // CK2: [[IFTHEN]]
-// CK2-DAG: call void @__tgt_target_data_end_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:.+]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+// CK2-DAG: call void @__tgt_target_exit_data_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:.+]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
 // CK2-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
 // CK2-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
 // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 // CK2-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0

 // RUN: %clang_cc1 -DCK3 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
 // SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
 #ifdef CK3
 
 // CK3-LABEL: no_target_devices
 void no_target_devices(int arg) {
-  // CK3-NOT: tgt_target_data_begin
+  // CK3-NOT: tgt_target_enter_data
   // CK3: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
-  // CK3-NOT: tgt_target_data_end
+  // CK3-NOT: tgt_target_exit_data
   // CK3: ret
   #pragma omp target exit data map(from: arg) if(arg) device(4)
   {++arg;}
 }
 #endif
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-64
 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s

 
 // CK4-LABEL: _Z3bari
 int bar(int arg){
   STT<int> A;
   return A.foo(arg);
 }
 
 // Region 00
-// CK4-NOT: __tgt_target_data_begin
+// CK4-NOT: __tgt_target_enter_data
 // CK4: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
 // CK4: [[IFTHEN]]
-// CK4-DAG: call void @__tgt_target_data_end_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:.+]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+// CK4-DAG: call void @__tgt_target_exit_data_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:.+]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
 // CK4-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
 // CK4-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
 // CK4-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK4-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 // CK4-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
 // CK4-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 // CK4-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0

clang/test/OpenMP/target_exit_data_depend_codegen.cpp

void foo(int arg) {
   // CK1: call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   // CK1: = call i32 [[TASK_ENTRY4]](i32 %{{.+}}, %struct.kmp_task_t_with_privates{{.+}}* [[RES_BC]])
   // CK1: call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   #pragma omp target exit data map(from:gb.b[:3]) depend(in: gb.b[:3], la, lb, gc, arg)
   {++arg;}
 }
 
 // CK1: define internal{{.*}} i32 [[TASK_ENTRY0]](i32{{.*}}, %struct.kmp_task_t_with_privates* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_end_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_exit_data_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** [[GEPM:%.+]])
 // CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
 // CK1-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 // CK1-DAG: [[GEPM]] = getelementptr inbounds {{.+}}[[M:%[^,]+]]
 // CK1-DAG: [[BP]] = load [1 x i8*]*, [1 x i8*]** [[BP_PRIV:%[^,]+]],
 // CK1-DAG: [[P]] = load [1 x i8*]*, [1 x i8*]** [[P_PRIV:%[^,]+]],
 // CK1-DAG: [[S]] = load [1 x i64]*, [1 x i64]** [[S_PRIV:%[^,]+]],
 // CK1-DAG: [[M]] = load [1 x i8*]*, [1 x i8*]** [[M_PRIV:%[^,]+]],
 // CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [1 x i8*]** [[BP_PRIV]], [1 x i8*]** [[P_PRIV]], [1 x i64]** [[S_PRIV]], [1 x i8*]** [[M_PRIV]])
 // CK1: ret i32 0
 // CK1: }
 
 // CK1: define internal{{.*}} i32 [[TASK_ENTRY2]](i32{{.*}}, %struct.kmp_task_t_with_privates{{.+}}* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_end_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** [[GEPM:%.+]])
 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 // CK1-DAG: [[GEPM]] = getelementptr inbounds {{.+}}[[M:%[^,]+]]
 // CK1-DAG: [[BP]] = load [1 x i8*]*, [1 x i8*]** [[BP_PRIV:%[^,]+]],
 // CK1-DAG: [[P]] = load [1 x i8*]*, [1 x i8*]** [[P_PRIV:%[^,]+]],
 // CK1-DAG: [[S]] = load [1 x i64]*, [1 x i64]** [[S_PRIV:%[^,]+]],
 // CK1-DAG: [[M]] = load [1 x i8*]*, [1 x i8*]** [[M_PRIV:%[^,]+]],
 // CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [1 x i8*]** [[BP_PRIV]], [1 x i8*]** [[P_PRIV]], [1 x i64]** [[S_PRIV]], [1 x i8*]** [[M_PRIV]])
 // CK1: ret i32 0
 // CK1: }
 
 // CK1: define internal{{.*}} i32 [[TASK_ENTRY3]](i32{{.*}}, %struct.kmp_task_t_with_privates{{.+}}* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** [[GEPM:%.+]])
 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 // CK1-DAG: [[GEPM]] = getelementptr inbounds {{.+}}[[M:%[^,]+]]
 // CK1-DAG: [[BP]] = load [1 x i8*]*, [1 x i8*]** [[BP_PRIV:%[^,]+]],
 // CK1-DAG: [[P]] = load [1 x i8*]*, [1 x i8*]** [[P_PRIV:%[^,]+]],
 // CK1-DAG: [[S]] = load [1 x i64]*, [1 x i64]** [[S_PRIV:%[^,]+]],
 // CK1-DAG: [[M]] = load [1 x i8*]*, [1 x i8*]** [[M_PRIV:%[^,]+]],
 // CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [1 x i8*]** [[BP_PRIV]], [1 x i8*]** [[P_PRIV]], [1 x i64]** [[S_PRIV]], [1 x i8*]** [[M_PRIV]])
-// CK1-NOT: __tgt_target_data_end_mapper
+// CK1-NOT: __tgt_target_exit_data_mapper
 // CK1: ret i32 0
 // CK1: }
 
 // CK1: define internal{{.*}} i32 [[TASK_ENTRY4]](i32{{.*}}, %struct.kmp_task_t_with_privates{{.+}}* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** [[GEPM:%.+]])
 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 // CK1-DAG: [[GEPM]] = getelementptr inbounds {{.+}}[[M:%[^,]+]]
 // CK1-DAG: [[BP]] = load [2 x i8*]*, [2 x i8*]** [[BP_PRIV:%[^,]+]],
 // CK1-DAG: [[P]] = load [2 x i8*]*, [2 x i8*]** [[P_PRIV:%[^,]+]],
 // CK1-DAG: [[S]] = load [2 x i64]*, [2 x i64]** [[S_PRIV:%[^,]+]],
 // CK1-DAG: [[M]] = load [2 x i8*]*, [2 x i8*]** [[M_PRIV:%[^,]+]],
 // CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [2 x i8*]** [[BP_PRIV]], [2 x i8*]** [[P_PRIV]], [2 x i64]** [[S_PRIV]], [2 x i8*]** [[M_PRIV]])
-// CK1-NOT: __tgt_target_data_end_mapper
+// CK1-NOT: __tgt_target_exit_data_mapper
 // CK1: ret i32 0
 // CK1: }
 
 #endif

clang/test/OpenMP/target_map_member_expr_array_section_codegen.cpp

maptest() {
     // CHECK: [[PTR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 1
     // CHECK: [[PTR1_DATA:%.+]] = bitcast i8** [[PTR1]] to float**
     // CHECK: store float* [[S_DATA_0_ADDR]], float** [[PTR1_DATA]],
     // CHECK: [[SIZE1:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 1
     // CHECK: store i64 24, i64* [[SIZE1]],
     // CHECK: [[BPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 0
     // CHECK: [[PTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 0
     // CHECK: [[SIZE:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 0
-    // CHECK: call void @__tgt_target_data_begin_mapper(i64 -1, i32 2, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAP_ENTER]], i32 0, i32 0), i8** null)
+    // CHECK: call void @__tgt_target_enter_data_mapper(i64 -1, i32 2, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAP_ENTER]], i32 0, i32 0), i8** null)
 #pragma omp target enter data map(alloc : s.data[:6])
   }
 
   ~maptest() {
     // CHECK: [[BPTRS:%.+]] = alloca [2 x i8*],
     // CHECK: [[PTRS:%.+]] = alloca [2 x i8*],
     // CHECK: [[SIZE:%.+]] = alloca [2 x i64],
     // CHECK: [[S_ADDR:%.+]] = getelementptr inbounds %struct.maptest, %struct.maptest* [[THIS:%.+]], i32 0, i32 0
~maptest() {
     // CHECK: [[PTR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 1
     // CHECK: [[PTR1_DATA:%.+]] = bitcast i8** [[PTR1]] to float**
     // CHECK: store float* [[S_DATA_0_ADDR]], float** [[PTR1_DATA]],
     // CHECK: [[SIZE1:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 1
     // CHECK: store i64 24, i64* [[SIZE1]],
     // CHECK: [[BPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 0
     // CHECK: [[PTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 0
     // CHECK: [[SIZE:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 0
-    // CHECK: call void @__tgt_target_data_end_mapper(i64 -1, i32 2, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAP_EXIT]], i32 0, i32 0), i8** null)
+    // CHECK: call void @__tgt_target_exit_data_mapper(i64 -1, i32 2, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAP_EXIT]], i32 0, i32 0), i8** null)
 #pragma omp target exit data map(delete : s.data[:6])
   }
 };
 
 maptest<float> a;
 
 #endif

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

 __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_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_enter_data_mapper, false, Void, Int64, Int32,
+          VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
+__OMP_RTL(__tgt_target_enter_data_nowait_mapper, false, Void, Int64, Int32,
+          VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
+__OMP_RTL(__tgt_target_exit_data_mapper, false, Void, Int64, Int32,
+          VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
+__OMP_RTL(__tgt_target_exit_data_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)
 __OMP_RTL(__tgt_mapper_num_components, false, Int64, VoidPtr)
 __OMP_RTL(__tgt_push_mapper_component, false, Void, VoidPtr, VoidPtr, VoidPtr,
           Int64, Int64)
 __OMP_RTL(__kmpc_task_allow_completion_event, false, VoidPtr, IdentPtr,

 __OMP_RTL_ATTRS(__tgt_target_teams_nowait_mapper, ForkAttrs, AttributeSet(), {})
 __OMP_RTL_ATTRS(__tgt_register_requires, ForkAttrs, AttributeSet(), {})
 __OMP_RTL_ATTRS(__tgt_target_data_begin_mapper, ForkAttrs, AttributeSet(), {})
 __OMP_RTL_ATTRS(__tgt_target_data_begin_nowait_mapper, ForkAttrs,
         AttributeSet(), {})
 __OMP_RTL_ATTRS(__tgt_target_data_end_mapper, ForkAttrs, AttributeSet(), {})
 __OMP_RTL_ATTRS(__tgt_target_data_end_nowait_mapper, ForkAttrs,
         AttributeSet(), {})
+__OMP_RTL_ATTRS(__tgt_target_enter_data_mapper, ForkAttrs, AttributeSet(), {})
+__OMP_RTL_ATTRS(__tgt_target_enter_data_nowait_mapper, ForkAttrs,
+        AttributeSet(), {})
+__OMP_RTL_ATTRS(__tgt_target_exit_data_mapper, ForkAttrs, AttributeSet(), {})
+__OMP_RTL_ATTRS(__tgt_target_exit_data_nowait_mapper, ForkAttrs,
+        AttributeSet(), {})
 __OMP_RTL_ATTRS(__tgt_target_data_update_mapper, ForkAttrs, AttributeSet(), {})
 __OMP_RTL_ATTRS(__tgt_target_data_update_nowait_mapper, ForkAttrs,
         AttributeSet(), {})
 __OMP_RTL_ATTRS(__tgt_mapper_num_components, ForkAttrs, AttributeSet(), {})
 __OMP_RTL_ATTRS(__tgt_push_mapper_component, ForkAttrs, AttributeSet(), {})
 __OMP_RTL_ATTRS(__kmpc_task_allow_completion_event, DefaultAttrs,
                 ReturnPtrAttrs, ParamAttrs(ReadOnlyPtrAttrs))
 

openmp/libomptarget/include/omptarget.h

void __tgt_target_data_end_mapper(int64_t device_id, int32_t arg_num,
                                   void **arg_mappers);
 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, int32_t depNum,
                                          void *depList, int32_t noAliasDepNum,
                                          void *noAliasDepList);
 
+void __tgt_target_enter_data_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);
+void __tgt_target_enter_data_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, int32_t depNum,
+    void *depList, int32_t noAliasDepNum, void *noAliasDepList);
+
+void __tgt_target_exit_data_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);
+void __tgt_target_exit_data_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, int32_t depNum,
+    void *depList, int32_t noAliasDepNum, void *noAliasDepList);
+
 /// passes data to/from the target
 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);
 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,
                                      int32_t depNum, void *depList,

openmp/libomptarget/src/exports

global:
     __tgt_target_teams;
     __tgt_target_data_begin_nowait;
     __tgt_target_data_end_nowait;
     __tgt_target_data_update_nowait;
     __tgt_target_nowait;
     __tgt_target_teams_nowait;
     __tgt_target_data_begin_mapper;
     __tgt_target_data_end_mapper;
+    __tgt_target_enter_data_mapper;
+    __tgt_target_exit_data_mapper;
     __tgt_target_data_update_mapper;
     __tgt_target_mapper;
     __tgt_target_teams_mapper;
     __tgt_target_data_begin_nowait_mapper;
     __tgt_target_data_end_nowait_mapper;
+    __tgt_target_enter_data_nowait_mapper;
+    __tgt_target_exit_data_nowait_mapper;
     __tgt_target_data_update_nowait_mapper;
     __tgt_target_nowait_mapper;
     __tgt_target_teams_nowait_mapper;
     __tgt_mapper_num_components;
     __tgt_push_mapper_component;
     omp_get_num_devices;
     omp_get_initial_device;
     omp_target_alloc;

openmp/libomptarget/src/interface.cpp

EXTERN void __tgt_target_data_end_nowait(int64_t device_id, int32_t arg_num,
     void *noAliasDepList) {
   if (depNum + noAliasDepNum > 0)
     __kmpc_omp_taskwait(NULL, __kmpc_global_thread_num(NULL));
 
   __tgt_target_data_end_mapper(device_id, arg_num, args_base, args, arg_sizes,
       arg_types, nullptr);
 }
 
-EXTERN 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) {
+static void 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, bool for_exit_data) {
   if (IsOffloadDisabled()) return;
   DP("Entering data end region with %d mappings\n", arg_num);
 
   // No devices available?
   if (device_id == OFFLOAD_DEVICE_DEFAULT) {
     device_id = omp_get_default_device();
   }
 
#ifdef OMPTARGET_DEBUG
   for (int i=0; i<arg_num; ++i) {
     DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
         ", Type=0x%" PRIx64 "\n", i, DPxPTR(args_base[i]), DPxPTR(args[i]),
         arg_sizes[i], arg_types[i]);
   }
 #endif
 
   int rc = target_data_end(Device, arg_num, args_base, args, arg_sizes,
-      arg_types, arg_mappers, nullptr);
+                           arg_types, arg_mappers, nullptr, for_exit_data);
   HandleTargetOutcome(rc == OFFLOAD_SUCCESS);
 }
 
+static void 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, int32_t depNum,
+    void *depList, int32_t noAliasDepNum, void *noAliasDepList,
+    bool for_exit_data) {
+  if (depNum + noAliasDepNum > 0)
+    __kmpc_omp_taskwait(NULL, __kmpc_global_thread_num(NULL));
+
+  target_data_end_mapper(device_id, arg_num, args_base, args, arg_sizes,
+                         arg_types, arg_mappers, for_exit_data);
+}
+
+EXTERN 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) {
+  target_data_end_mapper(device_id, arg_num, args_base, args, arg_sizes,
+                         arg_types, arg_mappers, /*for_exit_data=*/false);
+}
+
 EXTERN 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, int32_t depNum, void *depList,
     int32_t noAliasDepNum, void *noAliasDepList) {
-  if (depNum + noAliasDepNum > 0)
-    __kmpc_omp_taskwait(NULL, __kmpc_global_thread_num(NULL));
+  target_data_end_nowait_mapper(device_id, arg_num, args_base, args, arg_sizes,
+                                arg_types, arg_mappers, depNum, depList,
+                                noAliasDepNum, noAliasDepList,
+                                /*for_exit_data=*/false);
+}
 
-  __tgt_target_data_end_mapper(device_id, arg_num, args_base, args, arg_sizes,
-      arg_types, arg_mappers);
+EXTERN void __tgt_target_enter_data_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) {
+  return __tgt_target_data_begin_mapper(device_id, arg_num, args_base, args,
+                                        arg_sizes, arg_types, arg_mappers);
+}
+
+EXTERN void __tgt_target_enter_data_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, int32_t depNum,
+    void *depList, int32_t noAliasDepNum, void *noAliasDepList) {
+  return __tgt_target_data_begin_nowait_mapper(
+      device_id, arg_num, args_base, args, arg_sizes, arg_types, arg_mappers,
+      depNum, depList, noAliasDepNum, noAliasDepList);
+}
+
+EXTERN void __tgt_target_exit_data_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) {
+  return target_data_end_mapper(device_id, arg_num, args_base, args, arg_sizes,
+                                arg_types, arg_mappers, /*for_exit_data=*/true);
+}
+
+EXTERN void __tgt_target_exit_data_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, int32_t depNum,
+    void *depList, int32_t noAliasDepNum, void *noAliasDepList) {
+  return target_data_end_nowait_mapper(
+      device_id, arg_num, args_base, args, arg_sizes, arg_types, arg_mappers,
+      depNum, depList, noAliasDepNum, noAliasDepList, /*for_exit_data=*/true);
 }
 
 EXTERN 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) {
   __tgt_target_data_update_mapper(device_id, arg_num, args_base, args,
       arg_sizes, arg_types, nullptr);
 }
 

openmp/libomptarget/src/omptarget.cpp

if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) {
           Pointer_TgtPtrBegin, TgtPtrBase};
       Device.ShadowMtx.unlock();
     }
   }
 
   return OFFLOAD_SUCCESS;
 }
 
+static int target_data_end_not_for_exit_data(DeviceTy &Device, int32_t arg_num,
+                                             void **args_base, void **args,
+                                             int64_t *arg_sizes,
+                                             int64_t *arg_types,
+                                             void **arg_mappers,
+                                             __tgt_async_info *async_info_ptr) {
+  return target_data_end(Device, arg_num, args_base, args, arg_sizes, arg_types,
+                         arg_mappers, async_info_ptr, /*for_exit_data=*/false);
+}
+
+static int target_data_end_for_exit_data(DeviceTy &Device, int32_t arg_num,
+                                         void **args_base, void **args,
+                                         int64_t *arg_sizes, int64_t *arg_types,
+                                         void **arg_mappers,
+                                         __tgt_async_info *async_info_ptr) {
+  return target_data_end(Device, arg_num, args_base, args, arg_sizes, arg_types,
+                         arg_mappers, async_info_ptr, /*for_exit_data=*/true);
+}
+
 /// Internal function to undo the mapping and retrieve the data from the device.
 int target_data_end(DeviceTy &Device, int32_t arg_num, void **args_base,
                     void **args, int64_t *arg_sizes, int64_t *arg_types,
-                    void **arg_mappers, __tgt_async_info *async_info_ptr) {
+                    void **arg_mappers, __tgt_async_info *async_info_ptr,
+                    bool for_exit_data) {
   // process each input.
   for (int32_t i = arg_num - 1; i >= 0; --i) {
     // Ignore private variables and arrays - there is no mapping for them.
     // Also, ignore the use_device_ptr directive, it has no effect here.
     if ((arg_types[i] & OMP_TGT_MAPTYPE_LITERAL) ||
         (arg_types[i] & OMP_TGT_MAPTYPE_PRIVATE))
       continue;
 
     if (arg_mappers && arg_mappers[i]) {
       // Instead of executing the regular path of target_data_end, call the
       // target_data_mapper variant which will call target_data_end again
       // with new arguments.
       DP("Calling target_data_mapper for the %dth argument\n", i);
 
-      int rc = target_data_mapper(Device, args_base[i], args[i], arg_sizes[i],
-          arg_types[i], arg_mappers[i], target_data_end);
+      int rc =
+          target_data_mapper(Device, args_base[i], args[i], arg_sizes[i],
+                             arg_types[i], arg_mappers[i],
+                             for_exit_data ? target_data_end_for_exit_data
+                                           : target_data_end_not_for_exit_data);
 
       if (rc != OFFLOAD_SUCCESS) {
         DP("Call to target_data_end via target_data_mapper for custom mapper"
             " failed.\n");
         return OFFLOAD_FAIL;
       }
 
       // Skip the rest of this function, continue to the next argument.
for (int32_t i = arg_num - 1; i >= 0; --i) {
     bool HasPresentModifier = arg_types[i] & OMP_TGT_MAPTYPE_PRESENT;
 
     // If PTR_AND_OBJ, HstPtrBegin is address of pointee
     void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, data_size, IsLast,
         UpdateRef, IsHostPtr);
     if (!TgtPtrBegin && (data_size || HasPresentModifier)) {
       DP("Mapping does not exist (%s)\n",
          (HasPresentModifier ? "'present' map type modifier" : "ignored"));
-      if (HasPresentModifier) {
-        // FIXME: This should not be an error on exit from "omp target data",
-        // but it should be an error upon entering an "omp target exit data".
+      // If HasPresentModifier, complain if data is not present upon entering an
+      // "omp target exit data" but not upon exiting an "omp target data".
+      if (HasPresentModifier && for_exit_data) {
         MESSAGE("device mapping required by 'present' map type modifier does "
                 "not exist for host address " DPxMOD " (%ld bytes)",
                 DPxPTR(HstPtrBegin), data_size);
         return OFFLOAD_FAIL;
       }
     } else {
       DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
          " - is%s last\n",
for (auto it : fpArrays) {
     if (rt != OFFLOAD_SUCCESS) {
       DP("Deallocation of (first-)private arrays failed.\n");
       return OFFLOAD_FAIL;
     }
   }
 
   // Move data from device.
   int rt = target_data_end(Device, arg_num, args_base, args, arg_sizes,
-                           arg_types, arg_mappers, &AsyncInfo);
+                           arg_types, arg_mappers, &AsyncInfo,
+                           /*for_exit_data=*/false);
   if (rt != OFFLOAD_SUCCESS) {
     DP("Call to target_data_end failed, abort targe.\n");
     return OFFLOAD_FAIL;
   }
 
   if (Device.RTL->synchronize)
     return Device.RTL->synchronize(device_id, &AsyncInfo);
 
   return OFFLOAD_SUCCESS;
 }

openmp/libomptarget/src/private.h

 
 #ifndef _OMPTARGET_PRIVATE_H
 #define _OMPTARGET_PRIVATE_H
 
 #include <omptarget.h>
 
 #include <cstdint>
 
 extern int target_data_begin(DeviceTy &Device, int32_t arg_num,
                              void **args_base, void **args, int64_t *arg_sizes,
                              int64_t *arg_types, void **arg_mappers,
                              __tgt_async_info *async_info_ptr);
 
 extern int target_data_end(DeviceTy &Device, int32_t arg_num, void **args_base,
                            void **args, int64_t *arg_sizes, int64_t *arg_types,
-                           void **arg_mappers,
-                           __tgt_async_info *async_info_ptr);
+                           void **arg_mappers, __tgt_async_info *async_info_ptr,
+                           bool for_exit_data);
 
 extern int target_data_update(DeviceTy &Device, int32_t arg_num,
                               void **args_base, void **args,
                               int64_t *arg_sizes, int64_t *arg_types,
                               void **arg_mappers,
                               __tgt_async_info *async_info_ptr = nullptr);
 
 extern int target(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 team_num,
struct MapComponentInfoTy {
   MapComponentInfoTy(void *Base, void *Begin, int64_t Size, int64_t Type)
       : Base(Base), Begin(Begin), Size(Size), Type(Type) {}
 };
 
 // This structure stores all components of a user-defined mapper. The number of
 // components are dynamically decided, so we utilize C++ STL vector
 // implementation here.
 struct MapperComponentsTy {
   std::vector<MapComponentInfoTy> Components;
   int32_t size() { return Components.size(); }
 };
 
 // The mapper function pointer type. It follows the signature below:
 // void .omp_mapper.<type_name>.<mapper_id>.(void *rt_mapper_handle,
 //                                           void *base, void *begin,
 //                                           size_t size, int64_t type);
 typedef void (*MapperFuncPtrTy)(void *, void *, void *, int64_t, int64_t);
 
 // Function pointer type for target_data_* functions (target_data_begin,
 // target_data_end and target_data_update).
 typedef int (*TargetDataFuncPtrTy)(DeviceTy &, int32_t, void **, void **,
     int64_t *, int64_t *, void **, __tgt_async_info *);
 
 ////////////////////////////////////////////////////////////////////////////////
 // implementation for messages
 ////////////////////////////////////////////////////////////////////////////////
 
 #define MESSAGE0(_str)                                                         \
   do {                                                                         \

openmp/libomptarget/test/mapping/present/target_data_at_exit.c

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+#include <stdio.h>
+
+int main() {
+  int i;
+
+#pragma omp target enter data map(alloc:i)
+
+  // i isn't present at the end of the target data region, but the "present"
+  // modifier is only checked at the beginning of a region.
+#pragma omp target data map(present, alloc: i)
+  {
+#pragma omp target exit data map(delete:i)
+  }
+
+  // CHECK-NOT: Libomptarget
+  // CHECK: success
+  // CHECK-NOT: Libomptarget
+  fprintf(stderr, "success\n");
+
+  return 0;
+}