[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

8759

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

8912–8914

Do not append param here, use the one from Info

[jdenny]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8912–8914

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

public:
   virtual void emitNumTeamsClause(CodeGenFunction &CGF, const Expr *NumTeams,
                                   const Expr *ThreadLimit, SourceLocation Loc);
 
   /// Struct that keeps all the relevant information that should be kept
   /// throughout a 'target data' region.
   class TargetDataInfo {
     /// Set to true if device pointer information have to be obtained.
     bool RequiresDevicePointerInfo = false;
+    /// Set to true if Clang emits separate runtime calls for the beginning and
+    /// end of the region.  These calls might have separate map type arrays.
+    bool SeparateBeginEndCalls = false;
 
   public:
     /// The array of base pointer passed to the runtime library.
     llvm::Value *BasePointersArray = nullptr;
     /// The array of section pointers passed to the runtime library.
     llvm::Value *PointersArray = nullptr;
     /// The array of sizes passed to the runtime library.
     llvm::Value *SizesArray = nullptr;
-    /// The array of map types passed to the runtime library.
+    /// The array of map types passed to the runtime library for the beginning
+    /// of the region or for the entire region if there are no separate map
+    /// types for the region end.
     llvm::Value *MapTypesArray = nullptr;
+    /// The array of map types passed to the runtime library for the end of the
+    /// region, or nullptr if there are no separate map types for the region
+    /// end.
+    llvm::Value *MapTypesArrayEnd = nullptr;
     /// The array of user-defined mappers passed to the runtime library.
     llvm::Value *MappersArray = nullptr;
     /// Indicate whether any user-defined mapper exists.
     bool HasMapper = false;
     /// The total number of pointers passed to the runtime library.
     unsigned NumberOfPtrs = 0u;
     /// Map between the a declaration of a capture and the corresponding base
     /// pointer address where the runtime returns the device pointers.
     llvm::DenseMap<const ValueDecl *, Address> CaptureDeviceAddrMap;
 
     explicit TargetDataInfo() {}
-    explicit TargetDataInfo(bool RequiresDevicePointerInfo)
-        : RequiresDevicePointerInfo(RequiresDevicePointerInfo) {}
+    explicit TargetDataInfo(bool RequiresDevicePointerInfo,
+                            bool SeparateBeginEndCalls)
+        : RequiresDevicePointerInfo(RequiresDevicePointerInfo),
+          SeparateBeginEndCalls(SeparateBeginEndCalls) {}
     /// Clear information about the data arrays.
     void clearArrayInfo() {
       BasePointersArray = nullptr;
       PointersArray = nullptr;
       SizesArray = nullptr;
       MapTypesArray = nullptr;
+      MapTypesArrayEnd = nullptr;
       MappersArray = nullptr;
       HasMapper = false;
       NumberOfPtrs = 0u;
     }
     /// Return true if the current target data information has valid arrays.
     bool isValid() {
       return BasePointersArray && PointersArray && SizesArray &&
              MapTypesArray && (!HasMapper || MappersArray) && NumberOfPtrs;
     }
     bool requiresDevicePointerInfo() { return RequiresDevicePointerInfo; }
+    bool separateBeginEndCalls() { return SeparateBeginEndCalls; }
   };
 
   /// Emit the target data mapping code associated with \a D.
   /// \param D Directive to emit.
   /// \param IfCond Expression evaluated in if clause associated with the
   /// target directive, or null if no device clause is used.
   /// \param Device Expression evaluated in device clause associated with the
   /// target directive, or null if no device clause is used.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

 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
     // size so that a constant array could be eventually used.
     bool hasRuntimeEvaluationCaptureSize = false;
     for (llvm::Value *S : CombinedInfo.Sizes)
       if (!isa<llvm::Constant>(S)) {
std::string MaptypesName =
         CGM.getOpenMPRuntime().getName({"offload_maptypes"});
     auto *MapTypesArrayGbl = new llvm::GlobalVariable(
         CGM.getModule(), MapTypesArrayInit->getType(),
         /*isConstant=*/true, llvm::GlobalValue::PrivateLinkage,
         MapTypesArrayInit, MaptypesName);
     MapTypesArrayGbl->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
     Info.MapTypesArray = MapTypesArrayGbl;
 
+    // If there's a present map type modifier, it must not be applied to the end
+    // of a region, so generate a separate map type array in that case.
+    if (Info.separateBeginEndCalls()) {
+      bool EndMapTypesDiffer = false;
+      for (uint64_t &Type : Mapping) {
+        if (Type & MappableExprsHandler::OMP_MAP_PRESENT) {
+          Type &= ~MappableExprsHandler::OMP_MAP_PRESENT;
+          EndMapTypesDiffer = true;
+        }
+      }
+      if (EndMapTypesDiffer) {
+        MapTypesArrayInit =
+            llvm::ConstantDataArray::get(CGF.Builder.getContext(), Mapping);
+        MaptypesName = CGM.getOpenMPRuntime().getName({"offload_maptypes"});
+        MapTypesArrayGbl = new llvm::GlobalVariable(
+            CGM.getModule(), MapTypesArrayInit->getType(),
+            /*isConstant=*/true, llvm::GlobalValue::PrivateLinkage,
+            MapTypesArrayInit, MaptypesName);
+        MapTypesArrayGbl->setUnnamedAddr(
+            llvm::GlobalValue::UnnamedAddr::Global);
+        Info.MapTypesArrayEnd = MapTypesArrayGbl;
+      }
+    }
+
     for (unsigned I = 0; I < Info.NumberOfPtrs; ++I) {
       llvm::Value *BPVal = *CombinedInfo.BasePointers[I];
       llvm::Value *BP = CGF.Builder.CreateConstInBoundsGEP2_32(
           llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
           Info.BasePointersArray, 0, I);
       BP = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
           BP, BPVal->getType()->getPointerTo(/*AddrSpace=*/0));
       Address BPAddr(BP, Ctx.getTypeAlignInChars(Ctx.VoidPtrTy));
for (unsigned I = 0; I < Info.NumberOfPtrs; ++I) {
       }
       Address MAddr = CGF.Builder.CreateConstArrayGEP(MappersArray, I);
       CGF.Builder.CreateStore(MFunc, MAddr);
     }
   }
 }
 
 /// Emit the arguments to be passed to the runtime library based on the
-/// arrays of base pointers, pointers, sizes, map types, and mappers.
+/// arrays of base pointers, pointers, sizes, map types, and mappers.  If
+/// ForEndCall, emit map types to be passed for the end of the region instead of
+/// the beginning.
 static void emitOffloadingArraysArgument(
     CodeGenFunction &CGF, llvm::Value *&BasePointersArrayArg,
     llvm::Value *&PointersArrayArg, llvm::Value *&SizesArrayArg,
     llvm::Value *&MapTypesArrayArg, llvm::Value *&MappersArrayArg,
-    CGOpenMPRuntime::TargetDataInfo &Info) {
+    CGOpenMPRuntime::TargetDataInfo &Info, bool ForEndCall = false) {
+  assert((!ForEndCall || Info.separateBeginEndCalls()) &&
+         "expected region end call to runtime only when end call is separate");
   CodeGenModule &CGM = CGF.CGM;
   if (Info.NumberOfPtrs) {
     BasePointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
         llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
         Info.BasePointersArray,
         /*Idx0=*/0, /*Idx1=*/0);
     PointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
         llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
         Info.PointersArray,
         /*Idx0=*/0,
         /*Idx1=*/0);
     SizesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
         llvm::ArrayType::get(CGM.Int64Ty, Info.NumberOfPtrs), Info.SizesArray,
         /*Idx0=*/0, /*Idx1=*/0);
     MapTypesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32(
         llvm::ArrayType::get(CGM.Int64Ty, Info.NumberOfPtrs),
-        Info.MapTypesArray,
+        ForEndCall && Info.MapTypesArrayEnd ? Info.MapTypesArrayEnd
+                                            : Info.MapTypesArray,
         /*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);
auto &&BeginThenGen = [this, &D, Device, &Info,
 
     llvm::Value *BasePointersArrayArg = nullptr;
     llvm::Value *PointersArrayArg = nullptr;
     llvm::Value *SizesArrayArg = nullptr;
     llvm::Value *MapTypesArrayArg = nullptr;
     llvm::Value *MappersArrayArg = nullptr;
     emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
                                  SizesArrayArg, MapTypesArrayArg,
-                                 MappersArrayArg, Info);
+                                 MappersArrayArg, Info, /*ForEndCall=*/false);
 
     // Emit device ID if any.
     llvm::Value *DeviceID = nullptr;
     if (Device) {
       DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
                                            CGF.Int64Ty, /*isSigned=*/true);
     } else {
       DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);
auto &&EndThenGen = [this, Device, &Info](CodeGenFunction &CGF,
 
     llvm::Value *BasePointersArrayArg = nullptr;
     llvm::Value *PointersArrayArg = nullptr;
     llvm::Value *SizesArrayArg = nullptr;
     llvm::Value *MapTypesArrayArg = nullptr;
     llvm::Value *MappersArrayArg = nullptr;
     emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
                                  SizesArrayArg, MapTypesArrayArg,
-                                 MappersArrayArg, Info);
+                                 MappersArrayArg, Info, /*ForEndCall=*/true);
 
     // Emit device ID if any.
     llvm::Value *DeviceID = nullptr;
     if (Device) {
       DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
                                            CGF.Int64Ty, /*isSigned=*/true);
     } else {
       DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);

clang/lib/CodeGen/CGStmtOpenMP.cpp

for (const Expr *Ref : C.varlists()) {
 
     (void)PrivateScope.addPrivate(OrigVD, [PrivAddr]() { return PrivAddr; });
   }
 }
 
 // Generate the instructions for '#pragma omp target data' directive.
 void CodeGenFunction::EmitOMPTargetDataDirective(
     const OMPTargetDataDirective &S) {
-  CGOpenMPRuntime::TargetDataInfo Info(/*RequiresDevicePointerInfo=*/true);
+  CGOpenMPRuntime::TargetDataInfo Info(/*RequiresDevicePointerInfo=*/true,
+                                       /*SeparateBeginEndCalls=*/true);
 
   // Create a pre/post action to signal the privatization of the device pointer.
   // This action can be replaced by the OpenMP runtime code generation to
   // deactivate privatization.
   bool PrivatizeDevicePointers = false;
   class DevicePointerPrivActionTy : public PrePostActionTy {
     bool &PrivatizeDevicePointers;
 

clang/test/OpenMP/target_data_codegen.cpp

struct ST {
   T a;
   double *b;
 };
 
 ST<int> gb;
 double gc[100];
 
 // PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
-// CK1A: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+// CK1A: [[MTYPE00Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+
+// TARGET_PARAM=0x20 | TO=0x1 = 0x21
+// CK1A: [[MTYPE00End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x21]]]
 
 // PRESENT=0x1000 | CLOSE=0x400 | TARGET_PARAM=0x20 | ALWAYS=0x4 | TO=0x1 = 0x1425
-// CK1A: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1425]]]
+// CK1A: [[MTYPE01Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1425]]]
+
+// CLOSE=0x400 | TARGET_PARAM=0x20 | ALWAYS=0x4 | TO=0x1 = 0x425
+// CK1A: [[MTYPE01End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x425]]]
 
 // 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_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:32|64]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00Begin]]{{.+}})
   // 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-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK1A-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00End]]{{.+}})
   // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
   // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
   // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
   #pragma omp target data map(present, to: lb)
   {++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_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01Begin]]{{.+}})
   // 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-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+  // CK1A-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01End]]{{.+}})
   // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
   // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
   // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
   #pragma omp target data map(always close present, to: lb)
   {++arg;}
 
 }
 #endif

openmp/libomptarget/src/omptarget.cpp

for (int32_t I = ArgNum - 1; I >= 0; --I) {
 
     // If PTR_AND_OBJ, HstPtrBegin is address of pointee
     void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, DataSize, IsLast,
                                               UpdateRef, IsHostPtr);
     if (!TgtPtrBegin && (DataSize || 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".
+        // This should be an error upon entering an "omp target exit data".  It
+        // should not be an error upon exiting an "omp target data" or "omp
+        // target".  For "omp target data", Clang thus doesn't include present
+        // modifiers for end calls.  For "omp target", we have not found a valid
+        // OpenMP program for which the error matters: it appears that, if a
+        // program can guarantee that data is present at the beginning of an
+        // "omp target" region so that there's no error there, that data is also
+        // guaranteed to be present at the end.
         MESSAGE("device mapping required by 'present' map type modifier does "
                 "not exist for host address " DPxMOD " (%ld bytes)",
                 DPxPTR(HstPtrBegin), DataSize);
         return OFFLOAD_FAIL;
       }
     } else {
       DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
          " - is%s last\n",

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;
+}