[OpenMP] Change clang emitTargetDataCalls to use OMPIRBuilder

[TIFitis]

Summary

This patch changes the emitTargetDataCalls function in clang to make use of the OpenMPIRBuilder::createTargetData function for Target Data directive code gen.

Depends on D146557

[TIFitis]

clang/test/OpenMP/target_data_codegen.cpp

67–68

Similar to what I discussed for the if clause, the device clause argument need not be recomputed.
We can reuse the device argument from the begin mapper for the end mapper as well.

355–356

When both if clause and device clause are present, the device clause argument is inadvertently brought outside the IfThen region as we emit the llvm:Value from the Clang::Expr for it before making call to createTargetData.

I don't think this change would affect any use cases.

clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp

133–134

I think it is incorrect to recompute the branch condition here.

Consider the following C code:

int a = 1;
#pragma omp target data map(tofrom : a) if(a > 10) {a = 100;}

In this case the if condition is false before executing the region, but becomes true after.
If the branch condition is recomputed then it would take the IfThen branch here for executing the end_mapper call. This is incorrect and the end_mapper call should not be made partially without a begin _mapper call here.

Using the OMPIRBuilder fixes this.

[TIFitis]

Rebased

[TIFitis]

Ping for reviews :)

[jplehr]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10409

I think this can (and should be?) moved into the if in the following line as it is not used outside that scope

10413

capitalize here?

[TIFitis]

Addressed reviewer comments.

[jsjodin]

clang/test/OpenMP/target_data_codegen.cpp

355–356

When both if clause and device clause are present, the device clause argument is inadvertently brought outside the IfThen region as we emit the llvm:Value from the Clang::Expr for it before making call to createTargetData.

I don't think this change would affect any use cases.

Is it at all possible that the load could cause an exception if moved outside the if?

[TIFitis]

clang/test/OpenMP/target_data_codegen.cpp

355–356

Int pointer is not allowed inside device clause, so I don't think think the load can cause exceptions.

Also this would fix a scenario similar to the if clause, where in the following code target_begin would get device as 10 and target_end mapper call would get device as 100.

int arg = 10;
#pragma omp target data map(to: arg) if(arg < 20) device(arg)
  {arg = 100;}

[TIFitis]

Rebased

[TIFitis]

Rebased

[TIFitis]

Rebased

[jdoerfert]

This looks pretty good. The device vs if clause thing should be fixed while we are here.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10415

Move this behind the if generation and guard the emit scalar with a conditional, all is only needed if both Device is present and IfCond is not known true.

clang/test/OpenMP/target_data_codegen.cpp

355–356

Resting the value is good, computing it before the if is not. Take:

... if(isValidPtr(p)) device(*p)

We can generate the good code fairly easily, see above. Won't trigger much and will almost always be folded to straight line code anyway, except in cases like above, which is good. We end up with a phi and reuse the device value, which is also good. Win Win Win

clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp

133–134

Agreed. In general, we should always reuse the values we compute at the beginning of a region.

[jdoerfert]

clang/test/OpenMP/target_data_codegen.cpp

355–356

-Resting +Reusing (apply auto-correction)

[TIFitis]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10415

I have moved the device generation after the if generation.
But I am not sure what should be the guard here for emitting the scalar.

My idea was to generate something like:

%new_dev = alloca ...
if.then
    %x = load %device
    store %x %new_dev
    ...
    begin_mapper(..., %new_dev, ...)
...
if.then
    ...
    end_mapper(..., %new_dev, ...)

[TIFitis]

Moved device clause emition after if clause.

[TIFitis]

Gentle ping :)

[arsenm]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10403–10407

Why do you need to do this? I would expect EvaluateExprAsBool to handle the constant case for you

[TIFitis]

Addressed reviewer comments

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10403–10407

CGOpenMPRuntime::emitIfClause has the check for ConstantFoldsToSimpleInteger so I added it here.

Removing it didn't break any test so I've updated it,

[jdoerfert]

Let's move on. I described what this should look like, just for the record.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

10415

What we should do, though I don't think it will really matter, is the following:

device_val = undef;
if_val = eval_if_clause();
if (id_val) {
  device_val = eval_device_clause()
  ...other clauses ...
  begin_then(... Vals ...)
} else {
  begin_else(...)
}

if (if_val) {
  end_then(... Vals ...)
} else {
  end_else(...)
}

[TIFitis]

Rebased to fix build error.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

static void emitOffloadingArrays(
   Info.NumberOfPtrs = CombinedInfo.BasePointers.size();
 
   using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
   InsertPointTy AllocaIP(CGF.AllocaInsertPt->getParent(),
                          CGF.AllocaInsertPt->getIterator());
   InsertPointTy CodeGenIP(CGF.Builder.GetInsertBlock(),
                           CGF.Builder.GetInsertPoint());
 
-  auto fillInfoMap = [&](MappableExprsHandler::MappingExprInfo &MapExpr) {
+  auto FillInfoMap = [&](MappableExprsHandler::MappingExprInfo &MapExpr) {
     return emitMappingInformation(CGF, OMPBuilder, MapExpr);
   };
   if (CGM.getCodeGenOpts().getDebugInfo() !=
       llvm::codegenoptions::NoDebugInfo) {
     CombinedInfo.Names.resize(CombinedInfo.Exprs.size());
     llvm::transform(CombinedInfo.Exprs, CombinedInfo.Names.begin(),
-                    fillInfoMap);
+                    FillInfoMap);
   }
 
   auto DeviceAddrCB = [&](unsigned int I, llvm::Value *BP, llvm::Value *BPVal) {
     if (const ValueDecl *DevVD = CombinedInfo.DevicePtrDecls[I]) {
       Address BPAddr(BP, BPVal->getType(),
                      Ctx.getTypeAlignInChars(Ctx.VoidPtrTy));
       Info.CaptureDeviceAddrMap.try_emplace(DevVD, BPAddr);
     }
void CGOpenMPRuntime::emitTargetDataCalls(
     CGOpenMPRuntime::TargetDataInfo &Info) {
   if (!CGF.HaveInsertPoint())
     return;
 
   // Action used to replace the default codegen action and turn privatization
   // off.
   PrePostActionTy NoPrivAction;
 
-  // Generate the code for the opening of the data environment. Capture all the
-  // arguments of the runtime call by reference because they are used in the
-  // closing of the region.
-  auto &&BeginThenGen = [this, &D, Device, &Info,
-                         &CodeGen](CodeGenFunction &CGF, PrePostActionTy &) {
-    // Fill up the arrays with all the mapped variables.
-    MappableExprsHandler::MapCombinedInfoTy CombinedInfo;
-
-    // Get map clause information.
-    MappableExprsHandler MEHandler(D, CGF);
-    MEHandler.generateAllInfo(CombinedInfo);
-
-    // Fill up the arrays and create the arguments.
-    emitOffloadingArrays(CGF, CombinedInfo, Info, OMPBuilder,
-                         /*IsNonContiguous=*/true);
+  using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
+  InsertPointTy AllocaIP(CGF.AllocaInsertPt->getParent(),
+                         CGF.AllocaInsertPt->getIterator());
+  InsertPointTy CodeGenIP(CGF.Builder.GetInsertBlock(),
+                          CGF.Builder.GetInsertPoint());
+  llvm::OpenMPIRBuilder::LocationDescription OmpLoc(CodeGenIP);
 
-    llvm::OpenMPIRBuilder::TargetDataRTArgs RTArgs;
-    bool EmitDebug = CGF.CGM.getCodeGenOpts().getDebugInfo() !=
-                     llvm::codegenoptions::NoDebugInfo;
-    OMPBuilder.emitOffloadingArraysArgument(CGF.Builder, RTArgs, Info,
-                                            EmitDebug);
+  llvm::Value *IfCondVal = nullptr;
+  if (IfCond)
+    IfCondVal = CGF.EvaluateExprAsBool(IfCond);
 
-    // Emit device ID if any.
-    llvm::Value *DeviceID = nullptr;
-    if (Device) {
-      DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
-                                           CGF.Int64Ty, /*isSigned=*/true);
-    } else {
-      DeviceID = CGF.Builder.getInt64(OMP_DEVICEID_UNDEF);
-    }
 
-    // Emit the number of elements in the offloading arrays.
-    llvm::Value *PointerNum = CGF.Builder.getInt32(Info.NumberOfPtrs);
-    //
-    // Source location for the ident struct
-    llvm::Value *RTLoc = emitUpdateLocation(CGF, D.getBeginLoc());
-
-    llvm::Value *OffloadingArgs[] = {RTLoc,
-                                     DeviceID,
-                                     PointerNum,
-                                     RTArgs.BasePointersArray,
-                                     RTArgs.PointersArray,
-                                     RTArgs.SizesArray,
-                                     RTArgs.MapTypesArray,
-                                     RTArgs.MapNamesArray,
-                                     RTArgs.MappersArray};
-    CGF.EmitRuntimeCall(
-        OMPBuilder.getOrCreateRuntimeFunction(
-            CGM.getModule(), OMPRTL___tgt_target_data_begin_mapper),
-        OffloadingArgs);
+  // Fill up the arrays with all the mapped variables.
+  MappableExprsHandler::MapCombinedInfoTy CombinedInfo;
+  auto GenMapInfoCB =
+      [&](InsertPointTy CodeGenIP) -> llvm::OpenMPIRBuilder::MapInfosTy & {
+    CGF.Builder.restoreIP(CodeGenIP);
+    // Get map clause information.
+    MappableExprsHandler MEHandler(D, CGF);
+    MEHandler.generateAllInfo(CombinedInfo);
 
-    // If device pointer privatization is required, emit the body of the region
-    // here. It will have to be duplicated: with and without privatization.
-    if (!Info.CaptureDeviceAddrMap.empty())
-      CodeGen(CGF);
+    auto FillInfoMap = [&](MappableExprsHandler::MappingExprInfo &MapExpr) {
+      return emitMappingInformation(CGF, OMPBuilder, MapExpr);
     };
-
-  // Generate code for the closing of the data region.
-  auto &&EndThenGen = [this, Device, &Info, &D](CodeGenFunction &CGF,
-                                                PrePostActionTy &) {
-    assert(Info.isValid() && "Invalid data environment closing arguments.");
-
-    llvm::OpenMPIRBuilder::TargetDataRTArgs RTArgs;
-    bool EmitDebug = CGF.CGM.getCodeGenOpts().getDebugInfo() !=
-                     llvm::codegenoptions::NoDebugInfo;
-    OMPBuilder.emitOffloadingArraysArgument(CGF.Builder, RTArgs, Info,
-                                            EmitDebug,
-                                            /*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);
+    if (CGM.getCodeGenOpts().getDebugInfo() !=
+        llvm::codegenoptions::NoDebugInfo) {
+      CombinedInfo.Names.resize(CombinedInfo.Exprs.size());
+      llvm::transform(CombinedInfo.Exprs, CombinedInfo.Names.begin(),
+                      FillInfoMap);
     }
 
-    // Emit the number of elements in the offloading arrays.
-    llvm::Value *PointerNum = CGF.Builder.getInt32(Info.NumberOfPtrs);
-
-    // Source location for the ident struct
-    llvm::Value *RTLoc = emitUpdateLocation(CGF, D.getBeginLoc());
-
-    llvm::Value *OffloadingArgs[] = {RTLoc,
-                                     DeviceID,
-                                     PointerNum,
-                                     RTArgs.BasePointersArray,
-                                     RTArgs.PointersArray,
-                                     RTArgs.SizesArray,
-                                     RTArgs.MapTypesArray,
-                                     RTArgs.MapNamesArray,
-                                     RTArgs.MappersArray};
-    CGF.EmitRuntimeCall(
-        OMPBuilder.getOrCreateRuntimeFunction(
-            CGM.getModule(), OMPRTL___tgt_target_data_end_mapper),
-        OffloadingArgs);
+    return CombinedInfo;
   };
-
-  // If we need device pointer privatization, we need to emit the body of the
-  // region with no privatization in the 'else' branch of the conditional.
-  // Otherwise, we don't have to do anything.
-  auto &&BeginElseGen = [&Info, &CodeGen, &NoPrivAction](CodeGenFunction &CGF,
-                                                         PrePostActionTy &) {
+  using BodyGenTy = llvm::OpenMPIRBuilder::BodyGenTy;
+  auto BodyCB = [&](InsertPointTy CodeGenIP, BodyGenTy BodyGenType) {
+    CGF.Builder.restoreIP(CodeGenIP);
+    switch (BodyGenType) {
+    case BodyGenTy::Priv:
+      if (!Info.CaptureDeviceAddrMap.empty())
+        CodeGen(CGF);
+      break;
+    case BodyGenTy::DupNoPriv:
       if (!Info.CaptureDeviceAddrMap.empty()) {
         CodeGen.setAction(NoPrivAction);
         CodeGen(CGF);
       }
-  };
-
-  // We don't have to do anything to close the region if the if clause evaluates
-  // to false.
-  auto &&EndElseGen = [](CodeGenFunction &CGF, PrePostActionTy &) {};
-
-  if (IfCond) {
-    emitIfClause(CGF, IfCond, BeginThenGen, BeginElseGen);
-  } else {
-    RegionCodeGenTy RCG(BeginThenGen);
-    RCG(CGF);
-  }
-
-  // If we don't require privatization of device pointers, we emit the body in
-  // between the runtime calls. This avoids duplicating the body code.
+      break;
+    case BodyGenTy::NoPriv:
       if (Info.CaptureDeviceAddrMap.empty()) {
         CodeGen.setAction(NoPrivAction);
         CodeGen(CGF);
       }
+      break;
+    }
+    return InsertPointTy(CGF.Builder.GetInsertBlock(),
+                         CGF.Builder.GetInsertPoint());
+  };
 
-  if (IfCond) {
-    emitIfClause(CGF, IfCond, EndThenGen, EndElseGen);
-  } else {
-    RegionCodeGenTy RCG(EndThenGen);
-    RCG(CGF);
+  auto DeviceAddrCB = [&](unsigned int I, llvm::Value *BP, llvm::Value *BPVal) {
+    if (const ValueDecl *DevVD = CombinedInfo.DevicePtrDecls[I]) {
+      ASTContext &Ctx = CGF.getContext();
+      Address BPAddr(BP, BPVal->getType(),
+                     Ctx.getTypeAlignInChars(Ctx.VoidPtrTy));
+      Info.CaptureDeviceAddrMap.try_emplace(DevVD, BPAddr);
     }
+  };
+
+  auto CustomMapperCB = [&](unsigned int I) {
+    llvm::Value *MFunc = nullptr;
+    if (CombinedInfo.Mappers[I]) {
+      Info.HasMapper = true;
+      MFunc = CGF.CGM.getOpenMPRuntime().getOrCreateUserDefinedMapperFunc(
+          cast<OMPDeclareMapperDecl>(CombinedInfo.Mappers[I]));
+    }
+    return MFunc;
+  };
+
+  CGF.Builder.restoreIP(OMPBuilder.createTargetData(
+      OmpLoc, AllocaIP, CodeGenIP, DeviceID, IfCondVal, Info, GenMapInfoCB,
+      /*MapperFunc=*/nullptr, BodyCB, DeviceAddrCB, CustomMapperCB));
 }
 
 void CGOpenMPRuntime::emitTargetDataStandAloneCall(
     CodeGenFunction &CGF, const OMPExecutableDirective &D, const Expr *IfCond,
     const Expr *Device) {
   if (!CGF.HaveInsertPoint())
     return;
 

clang/test/OpenMP/target_data_codegen.cpp

void foo(int arg) {
 
   // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK1-DAG: store ptr @gc, ptr [[BP0]]
   // CK1-DAG: store ptr @gc, ptr [[P0]]
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 
-  // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV:%[^,]+]], i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[SIZE00]], ptr [[MTYPE00]], ptr null, ptr null)
-  // CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
-  // CK1-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
+  // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV]], i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[SIZE00]], ptr [[MTYPE00]], ptr null, ptr null)
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
   #pragma omp target data if(1+3-5) device(arg) map(from: gc)
   {++arg;}
 
   // Region 01
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   #pragma omp target data map(la) if(1+3-4)

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

 
 // CK2: [[IFELSE]]
 // CK2: br label %[[IFEND]]
 // CK2: [[IFEND]]
 // CK2: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 // CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
 
 // CK2: [[IFTHEN]]
-// CK2-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV:%[^,]+]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null)
-// CK2-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
-// CK2-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
+// CK2-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null)
 // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
 // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
 // CK2-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[PS]]
 // CK2: br label %[[IFEND:[^,]+]]
 // CK2: [[IFELSE]]
 // CK2: br label %[[IFEND]]
 // CK2: [[IFEND]]
 #endif

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

 
 // CK4: [[IFELSE]]
 // CK4: br label %[[IFEND]]
 // CK4: [[IFEND]]
 // CK4: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 // CK4: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
 
 // CK4: [[IFTHEN]]
-// CK4-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV:%[^,]+]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE00]], ptr null, ptr null)
-// CK4-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
-// CK4-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
+// CK4-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE00]], ptr null, ptr null)
 // CK4-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
 // CK4-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
 // CK4-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[PS]]
 // CK4: br label %[[IFEND:[^,]+]]
 // CK4: [[IFELSE]]
 // CK4: br label %[[IFEND]]
 // CK4: [[IFEND]]
 #endif

clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp

void foo(float *&lr, T *&tr) {
   // CK1:     [[BELSE]]:
   // CK1:     [[TTT:%.+]] = load ptr, ptr [[DECL]],
   // CK1:     getelementptr inbounds float, ptr [[TTT]], i32 1
   // CK1:     br label %[[BEND]]
   #pragma omp target data map(l[:10]) use_device_ptr(l) if(lr != 0)
   {
     ++l;
   }
   // CK1:     [[BEND]]:
-  // CK1:     [[CMP:%.+]] = icmp ne ptr %{{.+}}, null
   // CK1:     br i1 [[CMP]], label %[[BTHEN:.+]], label %[[BELSE:.+]]
 
   // CK1:     [[BTHEN]]:
   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE04]]
   // CK1:     br label %[[BEND:.+]]
 
   // CK1:     [[BELSE]]:
   // CK1:     br label %[[BEND]]

llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h

public:
   /// should be placed.
   /// \param IsBegin If true then emits begin mapper call otherwise emits
   /// end mapper call.
   /// \param DeviceID Stores the DeviceID from the device clause.
   /// \param IfCond Value which corresponds to the if clause condition.
   /// \param Info Stores all information realted to the Target Data directive.
   /// \param GenMapInfoCB Callback that populates the MapInfos and returns.
   /// \param BodyGenCB Optional Callback to generate the region code.
+  /// \param DeviceAddrCB Optional callback to generate code related to
+  /// use_device_ptr and use_device_addr.
+  /// \param CustomMapperCB Optional callback to generate code related to
+  /// custom mappers.
   OpenMPIRBuilder::InsertPointTy createTargetData(
       const LocationDescription &Loc, InsertPointTy AllocaIP,
       InsertPointTy CodeGenIP, Value *DeviceID, Value *IfCond,
       TargetDataInfo &Info,
       function_ref<MapInfosTy &(InsertPointTy CodeGenIP)> GenMapInfoCB,
       omp::RuntimeFunction *MapperFunc = nullptr,
       function_ref<InsertPointTy(InsertPointTy CodeGenIP,
                                  BodyGenTy BodyGenType)>
-          BodyGenCB = nullptr);
+          BodyGenCB = nullptr,
+      function_ref<void(unsigned int, Value *, Value *)> DeviceAddrCB = nullptr,
+      function_ref<Value *(unsigned int)> CustomMapperCB = nullptr);
 
   using TargetBodyGenCallbackTy = function_ref<InsertPointTy(
       InsertPointTy AllocaIP, InsertPointTy CodeGenIP)>;
 
   /// Generator for '#omp target'
   ///
   /// \param Loc where the target data construct was encountered.
   /// \param CodeGenIP The insertion point where the call to the outlined

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

 
 OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createTargetData(
     const LocationDescription &Loc, InsertPointTy AllocaIP,
     InsertPointTy CodeGenIP, Value *DeviceID, Value *IfCond,
     TargetDataInfo &Info,
     function_ref<MapInfosTy &(InsertPointTy CodeGenIP)> GenMapInfoCB,
     omp::RuntimeFunction *MapperFunc,
     function_ref<InsertPointTy(InsertPointTy CodeGenIP, BodyGenTy BodyGenType)>
-        BodyGenCB) {
+        BodyGenCB,
+    function_ref<void(unsigned int, Value *, Value *)> DeviceAddrCB,
+    function_ref<Value *(unsigned int)> CustomMapperCB) {
   if (!updateToLocation(Loc))
     return InsertPointTy();
 
   Builder.restoreIP(CodeGenIP);
   bool IsStandAlone = !BodyGenCB;
 
   // Generate the code for the opening of the data environment. Capture all the
   // arguments of the runtime call by reference because they are used in the
   // closing of the region.
   auto BeginThenGen = [&](InsertPointTy AllocaIP, InsertPointTy CodeGenIP) {
-    emitOffloadingArrays(AllocaIP, Builder.saveIP(),
-                         GenMapInfoCB(Builder.saveIP()), Info,
-                         /*IsNonContiguous=*/true);
+    emitOffloadingArrays(
+        AllocaIP, Builder.saveIP(), GenMapInfoCB(Builder.saveIP()), Info,
+        /*IsNonContiguous=*/true, DeviceAddrCB, CustomMapperCB);
 
     TargetDataRTArgs RTArgs;
     emitOffloadingArraysArgument(Builder, RTArgs, Info);
 
     // Emit the number of elements in the offloading arrays.
     Value *PointerNum = Builder.getInt32(Info.NumberOfPtrs);
 
     // Source location for the ident struct