[MLIR][OpenMP] Added MLIR translation support for use_device clauses

[TIFitis]

Summary

Added MLIR support for translating use_device_ptr and use_device_addr clauses for LLVMIR lowering.

• use_device_ptr: The mapped variables marked with use_device_ptr are accessed through a copy of the base pointer mappers. The mapper is copied onto a new temporary pointer variable.
• use_device_addr: The mapped variables marked with use_device_addr are accessed directly through the base pointer mappers.
• If mapping information is not provided explicitly then default map_type of alloc/release is assumed and the map_size is set to 0.

Depends on D152554

[TIFitis]

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

4215–4216

use_dev_ptr operands must always be pointer types so I think it is safe to always just create a ptr for them.

However if we want to support typed pointers then we would have to create something like:
Builder.CreateAlloca(UseDeviceOp.first->getType()->getPointerElementType())

This is not possible as getPointerElementType is deprecated and I couldn't find any alternative for this.

4335–4345

I have commented out some code in OpenMPToLLVMIRTranslation.cpp::1489. I think that is a better way of generating code for the use_dev_ptr, but the mapValue doesn't allow remapping.

This I am stuck with letting it use the old value inside the region and replacing it later here.

llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp

5081–5083

This unit test fails because of opaque pointer mismatch, The unit test runs with opaque pointers disabled however the my code only runs with it enabled.

I could not find a reasonable way to support typed pointers, I have explained this in another comment.

[TIFitis]

Split use_dev_ptr test into two.

[TIFitis]

Rebased

[TIFitis]

Need to change this patch to account for all the clang migration changes and make use of those functions.

[TIFitis]

Made changes to use updated OMPIRBuilder

[TIFitis]

Ping for review :)

[kiranchandramohan]

Could you describe the changes in the Summary? Particularly providing a short description of the code that is being emitted for use_device_ptr and use_device_addr?

[TIFitis]

Here is an example of what clang generates:

C code:

void foo(int *a, int *b, int *c){
#pragma omp target data map(from:b, c) use_device_ptr(a) use_device_addr(b){
    *a = 10;
    *b = 20;
    *c = 30;
  }
}

LLVM IR:

@.offload_sizes = private unnamed_addr constant [3 x i64] [i64 8, i64 8, i64 0]
@.offload_maptypes = private unnamed_addr constant [3 x i64] [i64 66, i64 2, i64 64]

define dso_local void @foo(ptr noundef %a, ptr noundef %b, ptr noundef %c) #0 {
entry:
  %a.addr = alloca ptr, align 8
  %b.addr = alloca ptr, align 8
  %c.addr = alloca ptr, align 8
  %.offload_baseptrs = alloca [3 x ptr], align 8
  %.offload_ptrs = alloca [3 x ptr], align 8
  %.offload_mappers = alloca [3 x ptr], align 8
  %0 = alloca ptr, align 8
  store ptr %a, ptr %a.addr, align 8
  store ptr %b, ptr %b.addr, align 8
  store ptr %c, ptr %c.addr, align 8
  %1 = load ptr, ptr %a.addr, align 8
  %2 = getelementptr inbounds [3 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
  store ptr %b.addr, ptr %2, align 8
  %3 = getelementptr inbounds [3 x ptr], ptr %.offload_ptrs, i32 0, i32 0
  store ptr %b.addr, ptr %3, align 8
  %4 = getelementptr inbounds [3 x ptr], ptr %.offload_mappers, i64 0, i64 0
  store ptr null, ptr %4, align 8
  %5 = getelementptr inbounds [3 x ptr], ptr %.offload_baseptrs, i32 0, i32 1
  store ptr %c.addr, ptr %5, align 8
  %6 = getelementptr inbounds [3 x ptr], ptr %.offload_ptrs, i32 0, i32 1
  store ptr %c.addr, ptr %6, align 8
  %7 = getelementptr inbounds [3 x ptr], ptr %.offload_mappers, i64 0, i64 1
  store ptr null, ptr %7, align 8
  %8 = getelementptr inbounds [3 x ptr], ptr %.offload_baseptrs, i32 0, i32 2
  store ptr %1, ptr %8, align 8
  %9 = getelementptr inbounds [3 x ptr], ptr %.offload_ptrs, i32 0, i32 2
  store ptr %1, ptr %9, align 8
  %10 = getelementptr inbounds [3 x ptr], ptr %.offload_mappers, i64 0, i64 2
  store ptr null, ptr %10, align 8
  %11 = getelementptr inbounds [3 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
  %12 = getelementptr inbounds [3 x ptr], ptr %.offload_ptrs, i32 0, i32 0
  call void @__tgt_target_data_begin_mapper(ptr @1, i64 -1, i32 3, ptr %11, ptr %12, ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
  %13 = load ptr, ptr %8, align 8
  store ptr %13, ptr %0, align 8
  %14 = load ptr, ptr %2, align 8
  %15 = load ptr, ptr %0, align 8
  store i32 10, ptr %15, align 4
  %16 = load ptr, ptr %14, align 8
  store i32 20, ptr %16, align 4
  %17 = load ptr, ptr %c.addr, align 8
  store i32 30, ptr %17, align 4
  %18 = getelementptr inbounds [3 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
  %19 = getelementptr inbounds [3 x ptr], ptr %.offload_ptrs, i32 0, i32 0
  call void @__tgt_target_data_end_mapper(ptr @1, i64 -1, i32 3, ptr %18, ptr %19, ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
  ret void
}

[TIFitis]

Ping for review

[kiranchandramohan]

I am starting to go through the patch.

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1462–1478

This block seems to be similar to devPtrOperands. Could you share code here?

mlir/test/Target/LLVMIR/omptarget-llvm.mlir

170–406

Please add a testcase with both use_device_ptr and use_device_addr.

[TIFitis]

Rebased. Addressed reviewer comments. Fixed bug in generating map info for device_addr - device info type was incorrectly being set to pointer.

[TIFitis]

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1462–1478

I've moved them into a functor.

[kiranchandramohan]

There is a build failure currently.

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1416

failure()?

1430

failure()?

[TIFitis]

Renamed failure() to fail()

[TIFitis]

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1416

I was calling the failed() functor defined above. I've renamed it to fail to avoid confusion.

[TIFitis]

Attempt at fixing build error

[kiranchandramohan]

While looking into this patch, I noticed that the current implementation in Semantics and MLIR requires the MAP clause to always be present. But the standard only says the following,

At least one map, use_device_addr or use_device_ptr clause must appear on the directive.

I will fix this. But this made having a look at the code generated only with use_device_addr, use_device_ptr and without map difficult.

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1408–1409

Nit:This only needs to take in as reference combinedInfo.

1557

Nit: expand the type here. Put an assert before this line for the op being a DataOp.

1561

Should this be an assert instead of an if? If not, Please add a comment on when it can and cannot be empty.

1563–1576

Might be worth creating a lambda and callling it separately with useDevPtrOperands, useDevAddrOperands.

1589

Should this be an assert instead of an if? If not, Please add a comment on when it can and cannot be empty.

mlir/test/Target/LLVMIR/omptarget-llvm.mlir

215

Is the default usage of dev_addr also with the pointer type?
If not using a single reference example might be easier to follow.

llvm.func @_QPopenmp_target_use_dev_addr() {
  %0 = llvm.mlir.constant(1 : i64) : i64
  %a = llvm.alloca %0 x !llvm.ptr<i32> : (i64) -> !llvm.ptr<i32>
  omp.target_data  map((from -> %a : !llvm.ptr<i32>)) use_device_addr(%a : !llvm.ptr<i32>)  {
  ^bb0(%arg0: !llvm.ptr<i32>): 
    %1 = llvm.mlir.constant(10 : i32) : i32
    llvm.store %1, %arg0 : !llvm.ptr<i32>
    omp.terminator
  }
  llvm.return
}

[TIFitis]

Addressed reviewer comments. Added missing LoadInst for use_device_addr inside region block.

[TIFitis]

In D146648#4527573, @kiranchandramohan wrote:

While looking into this patch, I noticed that the current implementation in Semantics and MLIR requires the MAP clause to always be present. But the standard only says the following,

At least one map, use_device_addr or use_device_ptr clause must appear on the directive.

I will fix this. But this made having a look at the code generated only with use_device_addr, use_device_ptr and without map difficult.

The test openmp_target_use_dev_addr_nomap does not provide any map clause for %a, so for %a marked with use_device_addr it has the same effect as no map clause being present. Key change when use_device_addr or use_device_ptr is used without providing any map information is that the map_size should be set to 0, and only the OMP_MAP_RETURN_PARAM flag should be set for the map_type.

[TIFitis]

BTW I am not sure how to resolve the windows build error that I am getting repeatedly. It's giving an unknown identifier error for the fail functor being called at OpenMPToLLVMIRTranslation.cpp:1429. Any insight would be much appreciated.

mlir/test/Target/LLVMIR/omptarget-llvm.mlir

215

I've added a openmp_target_use_dev_addr_no_ptr test

[kiranchandramohan]

BTW I am not sure how to resolve the windows build error that I am getting repeatedly. It's giving an unknown identifier error for the fail functor being called at OpenMPToLLVMIRTranslation.cpp:1429. Any insight would be much appreciated.

Might be a bug in the compiler in windows. You are trying to call a lambda function inside a lambda function here. Try moving the second lambda into a static function or try capturing the second lambda explicitly in the first lambda.

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1407–1409

Can you move this to just before the loop at 1406?

[TIFitis]

Addressed reviewer comments. Try to fix build error.

[TIFitis]

In D146648#4531472, @kiranchandramohan wrote:

BTW I am not sure how to resolve the windows build error that I am getting repeatedly. It's giving an unknown identifier error for the fail functor being called at OpenMPToLLVMIRTranslation.cpp:1429. Any insight would be much appreciated.

Might be a bug in the compiler in windows. You are trying to call a lambda function inside a lambda function here. Try moving the second lambda into a static function or try capturing the second lambda explicitly in the first lambda.

Thanks for the suggestion, explicitly capturing the lambda worked. The builds are all clean now.

[TIFitis]

Ping for review :)

[kiranchandramohan]

use_device_ptr: The mapped variables marked with use_device_ptr are accessed through new pointer allocas inside the region.
use_device_addr: The mapped variables marked with use_device_addr are accessed through the base pointer mappers.

This explanation in the summary talks of mapped variables. Mapping is always not necessary. Could you also expand the explanation a little more? And also talk a bit about the difference between use_device_ptr and use_device_addr.

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1562

If you do not need any codegen in this case for other switch options as well, then you can consider moving the if outside the switch.

1576–1577

Why is the load required here and not for devPtr?

[TIFitis]

In D146648#4554063, @kiranchandramohan wrote:

use_device_ptr: The mapped variables marked with use_device_ptr are accessed through new pointer allocas inside the region.
use_device_addr: The mapped variables marked with use_device_addr are accessed through the base pointer mappers.

This explanation in the summary talks of mapped variables. Mapping is always not necessary. Could you also expand the explanation a little more? And also talk a bit about the difference between use_device_ptr and use_device_addr.

The regular behaviour is for mapping to be present. In the event that mapping information is absent we simply don't have the map_type and size information available. As such according to the specification the size gets set to 0 and the map_type bits are omitted, i.e, default alloc/release is used as appropriate.

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

1562

No, we need the if condition here in combination with the switch case to generate code for device privatisation.

1576–1577

In case of dev_ptr the load inst is added by OMPIRBuilder itself along with a store Inst to store it into the new temp ptr var added.

For dev_addr we don't add the LoadInst in the OMPIRBuilder only because of compatibility issues with Clang.

[kiranchandramohan]

Looks OK. It will be good to have one more review.

[raghavendhra]

LGTM

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

function_ref<InsertPointTy(InsertPointTy CodeGenIP, BodyGenTy BodyGenType)>
         BodyGenCB,
     function_ref<void(unsigned int, Value *)> DeviceAddrCB,
     function_ref<Value *(unsigned int)> CustomMapperCB, Value *SrcLocInfo) {
   if (!updateToLocation(Loc))
     return InsertPointTy();
 
   Builder.restoreIP(CodeGenIP);
   bool IsStandAlone = !BodyGenCB;
   MapInfosTy *MapInfo;
   // 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) {
     MapInfo = &GenMapInfoCB(Builder.saveIP());
     emitOffloadingArrays(AllocaIP, Builder.saveIP(), *MapInfo, Info,
                          /*IsNonContiguous=*/true, DeviceAddrCB,
                          CustomMapperCB);
 
if (BodyGenCB) {
     } else {
       BeginThenGen(AllocaIP, Builder.saveIP());
     }
   }
 
   return Builder.saveIP();
 }
 
 // Copy input from pointer or i64 to the expected argument type.
 static Value *copyInput(IRBuilderBase &Builder, unsigned AddrSpace,
                         Value *Input, Argument &Arg) {
   auto Addr = Builder.CreateAlloca(Arg.getType()->isPointerTy()
                                        ? Arg.getType()
                                        : Type::getInt64Ty(Builder.getContext()),
                                    AddrSpace);
   auto AddrAscast =
       Builder.CreatePointerBitCastOrAddrSpaceCast(Addr, Input->getType());
   Builder.CreateStore(&Arg, AddrAscast);
   auto Copy = Builder.CreateLoad(Arg.getType(), AddrAscast);
 
   return Copy;
 }
 
 static Function *
 createOutlinedFunction(OpenMPIRBuilder &OMPBuilder, IRBuilderBase &Builder,
                        StringRef FuncName, SmallVectorImpl<Value *> &Inputs,
                        OpenMPIRBuilder::TargetBodyGenCallbackTy &CBFunc) {
Builder.CreateAlignedStore(
         BPVal, BP, M.getDataLayout().getPrefTypeAlign(PtrTy));
 
     if (Info.requiresDevicePointerInfo()) {
       if (CombinedInfo.DevicePointers[I] == DeviceInfoTy::Pointer) {
         CodeGenIP = Builder.saveIP();
         Builder.restoreIP(AllocaIP);
         Info.DevicePtrInfoMap[BPVal] = {BP, Builder.CreateAlloca(PtrTy)};
         Builder.restoreIP(CodeGenIP);
-        assert(DeviceAddrCB &&
-               "DeviceAddrCB missing for DevicePtr code generation");
+        if (DeviceAddrCB)
           DeviceAddrCB(I, Info.DevicePtrInfoMap[BPVal].second);
       } else if (CombinedInfo.DevicePointers[I] == DeviceInfoTy::Address) {
         Info.DevicePtrInfoMap[BPVal] = {BP, BP};
-        assert(DeviceAddrCB &&
-               "DeviceAddrCB missing for DevicePtr code generation");
+        if (DeviceAddrCB)
           DeviceAddrCB(I, BP);
       }
     }
 
     Value *PVal = CombinedInfo.Pointers[I];
     Value *P = Builder.CreateConstInBoundsGEP2_32(
         ArrayType::get(PtrTy, Info.NumberOfPtrs),
         Info.RTArgs.PointersArray, 0, I);
     // TODO: Check alignment correct.

llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp

TEST_F(OpenMPIRBuilderTest, TargetEnterData) {
   using InsertPointTy = OpenMPIRBuilder::InsertPointTy;
   auto GenMapInfoCB =
       [&](InsertPointTy codeGenIP) -> llvm::OpenMPIRBuilder::MapInfosTy & {
     // Get map clause information.
     Builder.restoreIP(codeGenIP);
 
     CombinedInfo.BasePointers.emplace_back(Val1);
     CombinedInfo.Pointers.emplace_back(Val1);
+    CombinedInfo.DevicePointers.emplace_back(
+        llvm::OpenMPIRBuilder::DeviceInfoTy::None);
     CombinedInfo.Sizes.emplace_back(Builder.getInt64(4));
     CombinedInfo.Types.emplace_back(llvm::omp::OpenMPOffloadMappingFlags(1));
     uint32_t temp;
     CombinedInfo.Names.emplace_back(
         OMPBuilder.getOrCreateSrcLocStr("unknown", temp));
     return CombinedInfo;
   };
 
TEST_F(OpenMPIRBuilderTest, TargetExitData) {
   using InsertPointTy = OpenMPIRBuilder::InsertPointTy;
   auto GenMapInfoCB =
       [&](InsertPointTy codeGenIP) -> llvm::OpenMPIRBuilder::MapInfosTy & {
     // Get map clause information.
     Builder.restoreIP(codeGenIP);
 
     CombinedInfo.BasePointers.emplace_back(Val1);
     CombinedInfo.Pointers.emplace_back(Val1);
+    CombinedInfo.DevicePointers.emplace_back(
+        llvm::OpenMPIRBuilder::DeviceInfoTy::None);
     CombinedInfo.Sizes.emplace_back(Builder.getInt64(4));
     CombinedInfo.Types.emplace_back(llvm::omp::OpenMPOffloadMappingFlags(2));
     uint32_t temp;
     CombinedInfo.Names.emplace_back(
         OMPBuilder.getOrCreateSrcLocStr("unknown", temp));
     return CombinedInfo;
   };
 
TEST_F(OpenMPIRBuilderTest, TargetDataRegion) {
   OpenMPIRBuilder::LocationDescription Loc({Builder.saveIP(), DL});
 
   int64_t DeviceID = 2;
 
   AllocaInst *Val1 =
       Builder.CreateAlloca(Builder.getInt32Ty(), Builder.getInt64(1));
   ASSERT_NE(Val1, nullptr);
 
+  AllocaInst *Val2 = Builder.CreateAlloca(Builder.getPtrTy());
+  ASSERT_NE(Val2, nullptr);
+
+  AllocaInst *Val3 = Builder.CreateAlloca(Builder.getPtrTy());
+  ASSERT_NE(Val3, nullptr);
+
   IRBuilder<>::InsertPoint AllocaIP(&F->getEntryBlock(),
                                     F->getEntryBlock().getFirstInsertionPt());
 
+  using DeviceInfoTy = llvm::OpenMPIRBuilder::DeviceInfoTy;
   llvm::OpenMPIRBuilder::MapInfosTy CombinedInfo;
   using InsertPointTy = OpenMPIRBuilder::InsertPointTy;
   auto GenMapInfoCB =
       [&](InsertPointTy codeGenIP) -> llvm::OpenMPIRBuilder::MapInfosTy & {
     // Get map clause information.
     Builder.restoreIP(codeGenIP);
+    uint32_t temp;
 
     CombinedInfo.BasePointers.emplace_back(Val1);
     CombinedInfo.Pointers.emplace_back(Val1);
+    CombinedInfo.DevicePointers.emplace_back(DeviceInfoTy::None);
     CombinedInfo.Sizes.emplace_back(Builder.getInt64(4));
     CombinedInfo.Types.emplace_back(llvm::omp::OpenMPOffloadMappingFlags(3));
-    uint32_t temp;
+    CombinedInfo.Names.emplace_back(
+        OMPBuilder.getOrCreateSrcLocStr("unknown", temp));
+
+    CombinedInfo.BasePointers.emplace_back(Val2);
+    CombinedInfo.Pointers.emplace_back(Val2);
+    CombinedInfo.DevicePointers.emplace_back(DeviceInfoTy::Pointer);
+    CombinedInfo.Sizes.emplace_back(Builder.getInt64(8));
+    CombinedInfo.Types.emplace_back(llvm::omp::OpenMPOffloadMappingFlags(67));
+    CombinedInfo.Names.emplace_back(
+        OMPBuilder.getOrCreateSrcLocStr("unknown", temp));
+
+    CombinedInfo.BasePointers.emplace_back(Val3);
+    CombinedInfo.Pointers.emplace_back(Val3);
+    CombinedInfo.DevicePointers.emplace_back(DeviceInfoTy::Address);
+    CombinedInfo.Sizes.emplace_back(Builder.getInt64(8));
+    CombinedInfo.Types.emplace_back(llvm::omp::OpenMPOffloadMappingFlags(67));
     CombinedInfo.Names.emplace_back(
         OMPBuilder.getOrCreateSrcLocStr("unknown", temp));
     return CombinedInfo;
   };
 
   llvm::OpenMPIRBuilder::TargetDataInfo Info(
-      /*RequiresDevicePointerInfo=*/false,
+      /*RequiresDevicePointerInfo=*/true,
       /*SeparateBeginEndCalls=*/true);
 
   OMPBuilder.Config.setIsGPU(true);
 
-  auto BodyCB = [&](InsertPointTy CodeGenIP, int BodyGenType) {
-    if (BodyGenType == 3) {
+  using BodyGenTy = llvm::OpenMPIRBuilder::BodyGenTy;
+  auto BodyCB = [&](InsertPointTy CodeGenIP, BodyGenTy BodyGenType) {
+    if (BodyGenType == BodyGenTy::Priv) {
+      EXPECT_EQ(Info.DevicePtrInfoMap.size(), 2u);
       Builder.restoreIP(CodeGenIP);
-      CallInst *TargetDataCall = dyn_cast<CallInst>(&BB->back());
+      CallInst *TargetDataCall =
+          dyn_cast<CallInst>(BB->back().getPrevNode()->getPrevNode());
       EXPECT_NE(TargetDataCall, nullptr);
       EXPECT_EQ(TargetDataCall->arg_size(), 9U);
       EXPECT_EQ(TargetDataCall->getCalledFunction()->getName(),
                 "__tgt_target_data_begin_mapper");
       EXPECT_TRUE(TargetDataCall->getOperand(1)->getType()->isIntegerTy(64));
       EXPECT_TRUE(TargetDataCall->getOperand(2)->getType()->isIntegerTy(32));
       EXPECT_TRUE(TargetDataCall->getOperand(8)->getType()->isPointerTy());
-      Builder.restoreIP(CodeGenIP);
+
+      LoadInst *LI = dyn_cast<LoadInst>(BB->back().getPrevNode());
+      EXPECT_NE(LI, nullptr);
+      StoreInst *SI = dyn_cast<StoreInst>(&BB->back());
+      EXPECT_NE(SI, nullptr);
+      EXPECT_EQ(SI->getValueOperand(), LI);
+      EXPECT_EQ(SI->getPointerOperand(), Info.DevicePtrInfoMap[Val2].second);
+      EXPECT_TRUE(isa<AllocaInst>(Info.DevicePtrInfoMap[Val2].second));
+      EXPECT_TRUE(isa<GetElementPtrInst>(Info.DevicePtrInfoMap[Val3].second));
       Builder.CreateStore(Builder.getInt32(99), Val1);
     }
     return Builder.saveIP();
   };
 
   Builder.restoreIP(OMPBuilder.createTargetData(
       Loc, AllocaIP, Builder.saveIP(), Builder.getInt64(DeviceID),
       /* IfCond= */ nullptr, Info, GenMapInfoCB, nullptr, BodyCB));
 
   CallInst *TargetDataCall = dyn_cast<CallInst>(&BB->back());
   EXPECT_NE(TargetDataCall, nullptr);
   EXPECT_EQ(TargetDataCall->arg_size(), 9U);
   EXPECT_EQ(TargetDataCall->getCalledFunction()->getName(),
             "__tgt_target_data_end_mapper");
   EXPECT_TRUE(TargetDataCall->getOperand(1)->getType()->isIntegerTy(64));
   EXPECT_TRUE(TargetDataCall->getOperand(2)->getType()->isIntegerTy(32));

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

convertOmpThreadprivate(Operation &opInst, llvm::IRBuilderBase &builder,
   llvm::Value *callInst =
       moduleTranslation.getOpenMPBuilder()->createCachedThreadPrivate(
           ompLoc, data, size, cacheName);
   llvm::Value *result = builder.CreateBitCast(callInst, globalValue->getType());
   moduleTranslation.mapValue(opInst.getResult(0), result);
   return success();
 }
 
-int64_t getSizeInBytes(DataLayout &DL, const mlir::Type &type) {
+int64_t getSizeInBytes(DataLayout &DL, const Type &type) {
   if (isa<LLVM::LLVMPointerType>(type))
     return DL.getTypeSize(cast<LLVM::LLVMPointerType>(type).getElementType());
 
   return 0;
 }
 
+// Generate all map related information and fill the combinedInfo.
 static void genMapInfos(llvm::IRBuilderBase &builder,
                         LLVM::ModuleTranslation &moduleTranslation,
                         DataLayout &DL,
                         llvm::OpenMPIRBuilder::MapInfosTy &combinedInfo,
                         const SmallVector<Value> &mapOperands,
-                        const ArrayAttr &mapTypes) {
-  // Get map clause information.
+                        const ArrayAttr &mapTypes,
+                        const SmallVector<Value> &devPtrOperands = {},
+                        const SmallVector<Value> &devAddrOperands = {}) {
   llvm::OpenMPIRBuilder *ompBuilder = moduleTranslation.getOpenMPBuilder();
-  unsigned index = 0;
-  for (const auto &mapOp : mapOperands) {
-    if (!mapOp.getType().isa<LLVM::LLVMPointerType>()) {
-      // TODO: Only LLVMPointerTypes are handled.
-      combinedInfo.BasePointers.clear();
-      combinedInfo.Pointers.clear();
-      combinedInfo.Sizes.clear();
-      combinedInfo.Types.clear();
-      combinedInfo.Names.clear();
-      return;
+
+  auto fail = [&combinedInfo]() -> void {
+    combinedInfo.DevicePointers.clear();
+  };
+
+  auto findMapInfo = [&combinedInfo](llvm::Value *val, unsigned &index) {
+    index = 0;
+    for (auto basePtr : combinedInfo.BasePointers) {
+      if (basePtr == val)
+        return true;
+      index++;
     }
+    return false;
+  };
+
+  unsigned index = 0;
+  for (const auto &mapOp : mapOperands) {
+    // TODO: Only LLVMPointerTypes are handled.
+    if (!mapOp.getType().isa<LLVM::LLVMPointerType>())
+      return fail();
 
     llvm::Value *mapOpValue = moduleTranslation.lookupValue(mapOp);
     combinedInfo.BasePointers.emplace_back(mapOpValue);
     combinedInfo.Pointers.emplace_back(mapOpValue);
+    combinedInfo.DevicePointers.emplace_back(
+        llvm::OpenMPIRBuilder::DeviceInfoTy::None);
     combinedInfo.Names.emplace_back(
-        mlir::LLVM::createMappingInformation(mapOp.getLoc(), *ompBuilder));
+        LLVM::createMappingInformation(mapOp.getLoc(), *ompBuilder));
     combinedInfo.Types.emplace_back(llvm::omp::OpenMPOffloadMappingFlags(
-        mapTypes[index].dyn_cast<mlir::IntegerAttr>().getInt()));
+        mapTypes[index].dyn_cast<IntegerAttr>().getInt()));
     combinedInfo.Sizes.emplace_back(
         builder.getInt64(getSizeInBytes(DL, mapOp.getType())));
     index++;
   }
+
+  auto addDevInfos = [&, fail](auto devOperands, auto devOpType) -> void {
+    for (const auto &devOp : devOperands) {
+      // TODO: Only LLVMPointerTypes are handled.
+      if (!devOp.getType().template isa<LLVM::LLVMPointerType>())
+        return fail();
+
+      llvm::Value *mapOpValue = moduleTranslation.lookupValue(devOp);
+
+      // Check if map info is already present for this entry.
+      unsigned infoIndex;
+      if (findMapInfo(mapOpValue, infoIndex)) {
+        combinedInfo.Types[infoIndex] |=
+            llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_RETURN_PARAM;
+        combinedInfo.DevicePointers[infoIndex] = devOpType;
+      } else {
+        combinedInfo.BasePointers.emplace_back(mapOpValue);
+        combinedInfo.Pointers.emplace_back(mapOpValue);
+        combinedInfo.DevicePointers.emplace_back(devOpType);
+        combinedInfo.Names.emplace_back(
+            LLVM::createMappingInformation(devOp.getLoc(), *ompBuilder));
+        combinedInfo.Types.emplace_back(
+            llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_RETURN_PARAM);
+        combinedInfo.Sizes.emplace_back(builder.getInt64(0));
+      }
+    }
+  };
+
+  addDevInfos(devPtrOperands, llvm::OpenMPIRBuilder::DeviceInfoTy::Pointer);
+  addDevInfos(devAddrOperands, llvm::OpenMPIRBuilder::DeviceInfoTy::Address);
 }
 
 static LogicalResult
 convertOmpTargetData(Operation *op, llvm::IRBuilderBase &builder,
                      LLVM::ModuleTranslation &moduleTranslation) {
   llvm::Value *ifCond = nullptr;
   int64_t deviceID = llvm::omp::OMP_DEVICEID_UNDEF;
   SmallVector<Value> mapOperands;
+  SmallVector<Value> useDevPtrOperands;
+  SmallVector<Value> useDevAddrOperands;
   ArrayAttr mapTypes;
   llvm::omp::RuntimeFunction RTLFn;
   DataLayout DL = DataLayout(op->getParentOfType<ModuleOp>());
 
   llvm::OpenMPIRBuilder *ompBuilder = moduleTranslation.getOpenMPBuilder();
 
   LogicalResult result =
       llvm::TypeSwitch<Operation *, LogicalResult>(op)
           .Case([&](omp::DataOp dataOp) {
-            if (!dataOp.getUseDeviceAddr().empty() ||
-                !dataOp.getUseDevicePtr().empty())
-              return failure();
-
             if (auto ifExprVar = dataOp.getIfExpr())
               ifCond = moduleTranslation.lookupValue(ifExprVar);
 
             if (auto devId = dataOp.getDevice())
-              if (auto constOp = mlir::dyn_cast<mlir::LLVM::ConstantOp>(
-                      devId.getDefiningOp()))
-                if (auto intAttr =
-                        dyn_cast<mlir::IntegerAttr>(constOp.getValue()))
+              if (auto constOp =
+                      dyn_cast<LLVM::ConstantOp>(devId.getDefiningOp()))
+                if (auto intAttr = dyn_cast<IntegerAttr>(constOp.getValue()))
                   deviceID = intAttr.getInt();
 
             mapOperands = dataOp.getMapOperands();
             if (dataOp.getMapTypes())
               mapTypes = dataOp.getMapTypes().value();
+            useDevPtrOperands = dataOp.getUseDevicePtr();
+            useDevAddrOperands = dataOp.getUseDeviceAddr();
             return success();
           })
           .Case([&](omp::EnterDataOp enterDataOp) {
             if (enterDataOp.getNowait())
               return failure();
 
             if (auto ifExprVar = enterDataOp.getIfExpr())
               ifCond = moduleTranslation.lookupValue(ifExprVar);
 
             if (auto devId = enterDataOp.getDevice())
-              if (auto constOp = mlir::dyn_cast<mlir::LLVM::ConstantOp>(
-                      devId.getDefiningOp()))
-                if (auto intAttr =
-                        dyn_cast<mlir::IntegerAttr>(constOp.getValue()))
+              if (auto constOp =
+                      dyn_cast<LLVM::ConstantOp>(devId.getDefiningOp()))
+                if (auto intAttr = dyn_cast<IntegerAttr>(constOp.getValue()))
                   deviceID = intAttr.getInt();
             RTLFn = llvm::omp::OMPRTL___tgt_target_data_begin_mapper;
             mapOperands = enterDataOp.getMapOperands();
             mapTypes = enterDataOp.getMapTypes();
             return success();
           })
           .Case([&](omp::ExitDataOp exitDataOp) {
             if (exitDataOp.getNowait())
               return failure();
 
             if (auto ifExprVar = exitDataOp.getIfExpr())
               ifCond = moduleTranslation.lookupValue(ifExprVar);
 
             if (auto devId = exitDataOp.getDevice())
-              if (auto constOp = mlir::dyn_cast<mlir::LLVM::ConstantOp>(
-                      devId.getDefiningOp()))
-                if (auto intAttr =
-                        dyn_cast<mlir::IntegerAttr>(constOp.getValue()))
+              if (auto constOp =
+                      dyn_cast<LLVM::ConstantOp>(devId.getDefiningOp()))
+                if (auto intAttr = dyn_cast<IntegerAttr>(constOp.getValue()))
                   deviceID = intAttr.getInt();
 
             RTLFn = llvm::omp::OMPRTL___tgt_target_data_end_mapper;
             mapOperands = exitDataOp.getMapOperands();
             mapTypes = exitDataOp.getMapTypes();
             return success();
           })
           .Default([&](Operation *op) {
             return op->emitError("unsupported OpenMP operation: ")
                    << op->getName();
           });
 
   if (failed(result))
     return failure();
 
   using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
 
   // Fill up the arrays with all the mapped variables.
   llvm::OpenMPIRBuilder::MapInfosTy combinedInfo;
   auto genMapInfoCB =
       [&](InsertPointTy codeGenIP) -> llvm::OpenMPIRBuilder::MapInfosTy & {
     builder.restoreIP(codeGenIP);
+    if (auto DataOp = dyn_cast<omp::DataOp>(op)) {
+      genMapInfos(builder, moduleTranslation, DL, combinedInfo, mapOperands,
+                  mapTypes, useDevPtrOperands, useDevAddrOperands);
+    } else {
       genMapInfos(builder, moduleTranslation, DL, combinedInfo, mapOperands,
                   mapTypes);
+    }
     return combinedInfo;
   };
 
-  LogicalResult bodyGenStatus = success();
+  llvm::OpenMPIRBuilder::TargetDataInfo info(/*RequiresDevicePointerInfo=*/true,
+                                             /*SeparateBeginEndCalls=*/true);
+
   using BodyGenTy = llvm::OpenMPIRBuilder::BodyGenTy;
+  LogicalResult bodyGenStatus = success();
   auto bodyGenCB = [&](InsertPointTy codeGenIP, BodyGenTy bodyGenType) {
+    assert(isa<omp::DataOp>(op) && "BodyGen requested for non DataOp");
+    Region &region = cast<omp::DataOp>(op).getRegion();
     switch (bodyGenType) {
     case BodyGenTy::Priv:
+      // Check if any device ptr/addr info is available
+      if (!info.DevicePtrInfoMap.empty()) {
+        builder.restoreIP(codeGenIP);
+        unsigned argIndex = 0;
+        for (auto &devPtrOp : useDevPtrOperands) {
+          llvm::Value *mapOpValue = moduleTranslation.lookupValue(devPtrOp);
+          const auto &arg = region.front().getArgument(argIndex);
+          moduleTranslation.mapValue(arg,
+                                     info.DevicePtrInfoMap[mapOpValue].second);
+          argIndex++;
+        }
+
+        for (auto &devAddrOp : useDevAddrOperands) {
+          llvm::Value *mapOpValue = moduleTranslation.lookupValue(devAddrOp);
+          const auto &arg = region.front().getArgument(argIndex);
+          auto *LI = builder.CreateLoad(
+              builder.getPtrTy(), info.DevicePtrInfoMap[mapOpValue].second);
+          moduleTranslation.mapValue(arg, LI);
+          argIndex++;
+        }
+
+        bodyGenStatus = inlineConvertOmpRegions(region, "omp.data.region",
+                                                builder, moduleTranslation);
+      }
       break;
     case BodyGenTy::DupNoPriv:
       break;
-    case BodyGenTy::NoPriv: {
-      // DataOp has only one region associated with it.
-      auto &region = cast<omp::DataOp>(op).getRegion();
+    case BodyGenTy::NoPriv:
+      // If device info is available then region has already been generated
+      if (info.DevicePtrInfoMap.empty()) {
         builder.restoreIP(codeGenIP);
         bodyGenStatus = inlineConvertOmpRegions(region, "omp.data.region",
                                                 builder, moduleTranslation);
       }
+      break;
     }
     return builder.saveIP();
   };
 
   llvm::OpenMPIRBuilder::LocationDescription ompLoc(builder);
   llvm::OpenMPIRBuilder::InsertPointTy allocaIP =
       findAllocaInsertPoint(builder, moduleTranslation);
-
-  // TODO: Add support for DevicePointerInfo
-  llvm::OpenMPIRBuilder::TargetDataInfo info(
-      /*RequiresDevicePointerInfo=*/false,
-      /*SeparateBeginEndCalls=*/true);
   if (isa<omp::DataOp>(op)) {
     builder.restoreIP(ompBuilder->createTargetData(
         ompLoc, allocaIP, builder.saveIP(), builder.getInt64(deviceID), ifCond,
         info, genMapInfoCB, nullptr, bodyGenCB));
   } else {
     builder.restoreIP(ompBuilder->createTargetData(
         ompLoc, allocaIP, builder.saveIP(), builder.getInt64(deviceID), ifCond,
         info, genMapInfoCB, &RTLFn));
convertDeclareTargetAttr(Operation *op,
   if (FunctionOpInterface funcOp = dyn_cast<FunctionOpInterface>(op)) {
     if (auto offloadMod = dyn_cast<omp::OffloadModuleInterface>(
             op->getParentOfType<ModuleOp>().getOperation())) {
       bool isDeviceCompilation = offloadMod.getIsTargetDevice();
       // FIXME: Temporarily disabled for host as it causes some issues when
       // lowering while removing functions at the current time.
       if (!isDeviceCompilation)
         return success();
-      
       omp::DeclareTargetDeviceType declareType =
           declareTargetAttr.getDeviceType().getValue();
 
       if ((isDeviceCompilation &&
            declareType == omp::DeclareTargetDeviceType::host) ||
           (!isDeviceCompilation &&
            declareType == omp::DeclareTargetDeviceType::nohost)) {
         llvm::Function *llvmFunc =

mlir/test/Target/LLVMIR/omptarget-llvm.mlir

 // CHECK:         %[[VAL_38:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_1]], i32 0, i32 0
 // CHECK:         call void @__tgt_target_data_end_mapper(ptr @3, i64 -1, i32 2, ptr %[[VAL_37]], ptr %[[VAL_38]], ptr @.offload_sizes.1, ptr @.offload_maptypes.2, ptr @.offload_mapnames.3, ptr null)
 // CHECK:         br label %[[VAL_39:.*]]
 // CHECK:       omp_if.else5:                                     ; preds = %[[VAL_25]]
 // CHECK:         br label %[[VAL_39]]
 // CHECK:       omp_if.end6:                                      ; preds = %[[VAL_30]], %[[VAL_29]]
 // CHECK:         ret void
 
 // -----
+
+llvm.func @_QPopenmp_target_use_dev_ptr() {
+  %0 = llvm.mlir.constant(1 : i64) : i64
+  %a = llvm.alloca %0 x !llvm.ptr<!llvm.ptr<i32>> : (i64) -> !llvm.ptr<!llvm.ptr<i32>>
+  omp.target_data  map((from -> %a : !llvm.ptr<!llvm.ptr<i32>>)) use_device_ptr(%a : !llvm.ptr<!llvm.ptr<i32>>)  {
+  ^bb0(%arg0: !llvm.ptr<!llvm.ptr<i32>>):
+    %1 = llvm.mlir.constant(10 : i32) : i32
+    %2 = llvm.load %arg0 : !llvm.ptr<!llvm.ptr<i32>>
+    llvm.store %1, %2 : !llvm.ptr<i32>
+    omp.terminator
+  }
+  llvm.return
+}
+
+// CHECK:         @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 8]
+// CHECK:         @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 66]
+// CHECK-LABEL: define void @_QPopenmp_target_use_dev_ptr
+// CHECK:         %[[VAL_0:.*]] = alloca [1 x ptr], align 8
+// CHECK:         %[[VAL_1:.*]] = alloca [1 x ptr], align 8
+// CHECK:         %[[VAL_2:.*]] = alloca [1 x ptr], align 8
+// CHECK:         %[[VAL_3:.*]] = alloca ptr, align 8
+// CHECK:         %[[VAL_4:.*]] = alloca ptr, i64 1, align 8
+// CHECK:         br label %[[VAL_5:.*]]
+// CHECK:       entry:                                            ; preds = %[[VAL_6:.*]]
+// CHECK:         %[[VAL_7:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         store ptr %[[VAL_4]], ptr %[[VAL_7]], align 8
+// CHECK:         %[[VAL_8:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         store ptr %[[VAL_4]], ptr %[[VAL_8]], align 8
+// CHECK:         %[[VAL_9:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_2]], i64 0, i64 0
+// CHECK:         store ptr null, ptr %[[VAL_9]], align 8
+// CHECK:         %[[VAL_10:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         %[[VAL_11:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         call void @__tgt_target_data_begin_mapper(ptr @2, i64 -1, i32 1, ptr %[[VAL_10]], ptr %[[VAL_11]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
+// CHECK:         %[[VAL_12:.*]] = load ptr, ptr %[[VAL_7]], align 8
+// CHECK:         store ptr %[[VAL_12]], ptr %[[VAL_3]], align 8
+// CHECK:         %[[VAL_13:.*]] = load ptr, ptr %[[VAL_3]], align 8
+// CHECK:         store i32 10, ptr %[[VAL_13]], align 4
+// CHECK:         %[[VAL_14:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         %[[VAL_15:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         call void @__tgt_target_data_end_mapper(ptr @2, i64 -1, i32 1, ptr %[[VAL_14]], ptr %[[VAL_15]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
+// CHECK:         ret void
+
+// -----
+
+llvm.func @_QPopenmp_target_use_dev_addr() {
+  %0 = llvm.mlir.constant(1 : i64) : i64
+  %a = llvm.alloca %0 x !llvm.ptr<!llvm.ptr<i32>> : (i64) -> !llvm.ptr<!llvm.ptr<i32>>
+  omp.target_data  map((from -> %a : !llvm.ptr<!llvm.ptr<i32>>)) use_device_addr(%a : !llvm.ptr<!llvm.ptr<i32>>)  {
+  ^bb0(%arg0: !llvm.ptr<!llvm.ptr<i32>>):
+    %1 = llvm.mlir.constant(10 : i32) : i32
+    %2 = llvm.load %arg0 : !llvm.ptr<!llvm.ptr<i32>>
+    llvm.store %1, %2 : !llvm.ptr<i32>
+    omp.terminator
+  }
+  llvm.return
+}
+
+// CHECK:         @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 8]
+// CHECK:         @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 66]
+// CHECK-LABEL: define void @_QPopenmp_target_use_dev_addr
+// CHECK:         %[[VAL_0:.*]] = alloca [1 x ptr], align 8
+// CHECK:         %[[VAL_1:.*]] = alloca [1 x ptr], align 8
+// CHECK:         %[[VAL_2:.*]] = alloca [1 x ptr], align 8
+// CHECK:         %[[VAL_3:.*]] = alloca ptr, i64 1, align 8
+// CHECK:         br label %[[VAL_4:.*]]
+// CHECK:       entry:                                            ; preds = %[[VAL_5:.*]]
+// CHECK:         %[[VAL_6:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         store ptr %[[VAL_3]], ptr %[[VAL_6]], align 8
+// CHECK:         %[[VAL_7:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         store ptr %[[VAL_3]], ptr %[[VAL_7]], align 8
+// CHECK:         %[[VAL_8:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_2]], i64 0, i64 0
+// CHECK:         store ptr null, ptr %[[VAL_8]], align 8
+// CHECK:         %[[VAL_9:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         %[[VAL_10:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         call void @__tgt_target_data_begin_mapper(ptr @2, i64 -1, i32 1, ptr %[[VAL_9]], ptr %[[VAL_10]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
+// CHECK:         %[[VAL_11:.*]] = load ptr, ptr %[[VAL_6]], align 8
+// CHECK:         %[[VAL_12:.*]] = load ptr, ptr %[[VAL_11]], align 8
+// CHECK:         store i32 10, ptr %[[VAL_12]], align 4
+// CHECK:         %[[VAL_13:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         %[[VAL_14:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         call void @__tgt_target_data_end_mapper(ptr @2, i64 -1, i32 1, ptr %[[VAL_13]], ptr %[[VAL_14]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
+// CHECK:         ret void
+
+// -----
+
+llvm.func @_QPopenmp_target_use_dev_addr_no_ptr() {
+  %0 = llvm.mlir.constant(1 : i64) : i64
+  %a = llvm.alloca %0 x !llvm.ptr<i32> : (i64) -> !llvm.ptr<i32>
+  omp.target_data  map((tofrom -> %a : !llvm.ptr<i32>)) use_device_addr(%a : !llvm.ptr<i32>)  {
+  ^bb0(%arg0: !llvm.ptr<i32>):
+    %1 = llvm.mlir.constant(10 : i32) : i32
+    llvm.store %1, %arg0 : !llvm.ptr<i32>
+    omp.terminator
+  }
+  llvm.return
+}
+
+// CHECK:         @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 4]
+// CHECK:         @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 67]
+// CHECK-LABEL: define void @_QPopenmp_target_use_dev_addr_no_ptr
+// CHECK:         %[[VAL_0:.*]] = alloca [1 x ptr], align 8
+// CHECK:         %[[VAL_1:.*]] = alloca [1 x ptr], align 8
+// CHECK:         %[[VAL_2:.*]] = alloca [1 x ptr], align 8
+// CHECK:         %[[VAL_3:.*]] = alloca i32, i64 1, align 4
+// CHECK:         br label %[[VAL_4:.*]]
+// CHECK:       entry:                                            ; preds = %[[VAL_5:.*]]
+// CHECK:         %[[VAL_6:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         store ptr %[[VAL_3]], ptr %[[VAL_6]], align 8
+// CHECK:         %[[VAL_7:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         store ptr %[[VAL_3]], ptr %[[VAL_7]], align 8
+// CHECK:         %[[VAL_8:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_2]], i64 0, i64 0
+// CHECK:         store ptr null, ptr %[[VAL_8]], align 8
+// CHECK:         %[[VAL_9:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         %[[VAL_10:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         call void @__tgt_target_data_begin_mapper(ptr @2, i64 -1, i32 1, ptr %[[VAL_9]], ptr %[[VAL_10]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
+// CHECK:         %[[VAL_11:.*]] = load ptr, ptr %[[VAL_6]], align 8
+// CHECK:         store i32 10, ptr %[[VAL_11]], align 4
+// CHECK:         %[[VAL_12:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         %[[VAL_13:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         call void @__tgt_target_data_end_mapper(ptr @2, i64 -1, i32 1, ptr %[[VAL_12]], ptr %[[VAL_13]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
+// CHECK:         ret void
+
+// -----
+
+llvm.func @_QPopenmp_target_use_dev_addr_nomap() {
+  %0 = llvm.mlir.constant(1 : i64) : i64
+  %a = llvm.alloca %0 x !llvm.ptr<!llvm.ptr<i32>> : (i64) -> !llvm.ptr<!llvm.ptr<i32>>
+  %1 = llvm.mlir.constant(1 : i64) : i64
+  %b = llvm.alloca %0 x !llvm.ptr<!llvm.ptr<i32>> : (i64) -> !llvm.ptr<!llvm.ptr<i32>>
+  omp.target_data  map((from -> %b : !llvm.ptr<!llvm.ptr<i32>>)) use_device_addr(%a : !llvm.ptr<!llvm.ptr<i32>>)  {
+  ^bb0(%arg0: !llvm.ptr<!llvm.ptr<i32>>):
+    %2 = llvm.mlir.constant(10 : i32) : i32
+    %3 = llvm.load %arg0 : !llvm.ptr<!llvm.ptr<i32>>
+    llvm.store %2, %3 : !llvm.ptr<i32>
+    %4 = llvm.mlir.constant(20 : i32) : i32
+    %5 = llvm.load %b : !llvm.ptr<!llvm.ptr<i32>>
+    llvm.store %4, %5 : !llvm.ptr<i32>
+    omp.terminator
+  }
+  llvm.return
+}
+
+// CHECK:         @.offload_sizes = private unnamed_addr constant [2 x i64] [i64 8, i64 0]
+// CHECK:         @.offload_maptypes = private unnamed_addr constant [2 x i64] [i64 2, i64 64]
+// CHECK-LABEL: define void @_QPopenmp_target_use_dev_addr_nomap
+// CHECK:         %[[VAL_0:.*]] = alloca [2 x ptr], align 8
+// CHECK:         %[[VAL_1:.*]] = alloca [2 x ptr], align 8
+// CHECK:         %[[VAL_2:.*]] = alloca [2 x ptr], align 8
+// CHECK:         %[[VAL_3:.*]] = alloca ptr, i64 1, align 8
+// CHECK:         %[[VAL_4:.*]] = alloca ptr, i64 1, align 8
+// CHECK:         br label %[[VAL_5:.*]]
+// CHECK:       entry:                                            ; preds = %[[VAL_6:.*]]
+// CHECK:         %[[VAL_7:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         store ptr %[[VAL_4]], ptr %[[VAL_7]], align 8
+// CHECK:         %[[VAL_8:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         store ptr %[[VAL_4]], ptr %[[VAL_8]], align 8
+// CHECK:         %[[VAL_9:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_2]], i64 0, i64 0
+// CHECK:         store ptr null, ptr %[[VAL_9]], align 8
+// CHECK:         %[[VAL_10:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_0]], i32 0, i32 1
+// CHECK:         store ptr %[[VAL_3]], ptr %[[VAL_10]], align 8
+// CHECK:         %[[VAL_11:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_1]], i32 0, i32 1
+// CHECK:         store ptr %[[VAL_3]], ptr %[[VAL_11]], align 8
+// CHECK:         %[[VAL_12:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_2]], i64 0, i64 1
+// CHECK:         store ptr null, ptr %[[VAL_12]], align 8
+// CHECK:         %[[VAL_13:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         %[[VAL_14:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         call void @__tgt_target_data_begin_mapper(ptr @3, i64 -1, i32 2, ptr %[[VAL_13]], ptr %[[VAL_14]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
+// CHECK:         %[[VAL_15:.*]] = load ptr, ptr %[[VAL_10]], align 8
+// CHECK:         %[[VAL_16:.*]] = load ptr, ptr %[[VAL_15]], align 8
+// CHECK:         store i32 10, ptr %[[VAL_16]], align 4
+// CHECK:         %[[VAL_17:.*]] = load ptr, ptr %[[VAL_4]], align 8
+// CHECK:         store i32 20, ptr %[[VAL_17]], align 4
+// CHECK:         %[[VAL_18:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         %[[VAL_19:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         call void @__tgt_target_data_end_mapper(ptr @3, i64 -1, i32 2, ptr %[[VAL_18]], ptr %[[VAL_19]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
+// CHECK:         ret void
+
+// -----
+
+llvm.func @_QPopenmp_target_use_dev_both() {
+  %0 = llvm.mlir.constant(1 : i64) : i64
+  %a = llvm.alloca %0 x !llvm.ptr<!llvm.ptr<i32>> : (i64) -> !llvm.ptr<!llvm.ptr<i32>>
+  %1 = llvm.mlir.constant(1 : i64) : i64
+  %b = llvm.alloca %0 x !llvm.ptr<!llvm.ptr<i32>> : (i64) -> !llvm.ptr<!llvm.ptr<i32>>
+  omp.target_data  map((tofrom -> %a : !llvm.ptr<!llvm.ptr<i32>>), (tofrom -> %b : !llvm.ptr<!llvm.ptr<i32>>)) use_device_ptr(%a : !llvm.ptr<!llvm.ptr<i32>>) use_device_addr(%b : !llvm.ptr<!llvm.ptr<i32>>)  {
+  ^bb0(%arg0: !llvm.ptr<!llvm.ptr<i32>>, %arg1: !llvm.ptr<!llvm.ptr<i32>>):
+    %2 = llvm.mlir.constant(10 : i32) : i32
+    %3 = llvm.load %arg0 : !llvm.ptr<!llvm.ptr<i32>>
+    llvm.store %2, %3 : !llvm.ptr<i32>
+    %4 = llvm.mlir.constant(20 : i32) : i32
+    %5 = llvm.load %arg1 : !llvm.ptr<!llvm.ptr<i32>>
+    llvm.store %4, %5 : !llvm.ptr<i32>
+    omp.terminator
+  }
+  llvm.return
+}
+
+// CHECK:         @.offload_sizes = private unnamed_addr constant [2 x i64] [i64 8, i64 8]
+// CHECK:         @.offload_maptypes = private unnamed_addr constant [2 x i64] [i64 67, i64 67]
+// CHECK-LABEL: define void @_QPopenmp_target_use_dev_both
+// CHECK:         %[[VAL_0:.*]] = alloca [2 x ptr], align 8
+// CHECK:         %[[VAL_1:.*]] = alloca [2 x ptr], align 8
+// CHECK:         %[[VAL_2:.*]] = alloca [2 x ptr], align 8
+// CHECK:         %[[VAL_3:.*]] = alloca ptr, align 8
+// CHECK:         %[[VAL_4:.*]] = alloca ptr, i64 1, align 8
+// CHECK:         %[[VAL_5:.*]] = alloca ptr, i64 1, align 8
+// CHECK:         br label %[[VAL_6:.*]]
+// CHECK:       entry:                                            ; preds = %[[VAL_7:.*]]
+// CHECK:         %[[VAL_8:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         store ptr %[[VAL_4]], ptr %[[VAL_8]], align 8
+// CHECK:         %[[VAL_9:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         store ptr %[[VAL_4]], ptr %[[VAL_9]], align 8
+// CHECK:         %[[VAL_10:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_2]], i64 0, i64 0
+// CHECK:         store ptr null, ptr %[[VAL_10]], align 8
+// CHECK:         %[[VAL_11:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_0]], i32 0, i32 1
+// CHECK:         store ptr %[[VAL_5]], ptr %[[VAL_11]], align 8
+// CHECK:         %[[VAL_12:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_1]], i32 0, i32 1
+// CHECK:         store ptr %[[VAL_5]], ptr %[[VAL_12]], align 8
+// CHECK:         %[[VAL_13:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_2]], i64 0, i64 1
+// CHECK:         store ptr null, ptr %[[VAL_13]], align 8
+// CHECK:         %[[VAL_14:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         %[[VAL_15:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         call void @__tgt_target_data_begin_mapper(ptr @3, i64 -1, i32 2, ptr %[[VAL_14]], ptr %[[VAL_15]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
+// CHECK:         %[[VAL_16:.*]] = load ptr, ptr %[[VAL_8]], align 8
+// CHECK:         store ptr %[[VAL_16]], ptr %[[VAL_3]], align 8
+// CHECK:         %[[VAL_17:.*]] = load ptr, ptr %[[VAL_11]], align 8
+// CHECK:         %[[VAL_18:.*]] = load ptr, ptr %[[VAL_3]], align 8
+// CHECK:         store i32 10, ptr %[[VAL_18]], align 4
+// CHECK:         %[[VAL_19:.*]] = load ptr, ptr %[[VAL_17]], align 8
+// CHECK:         store i32 20, ptr %[[VAL_19]], align 4
+// CHECK:         %[[VAL_20:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_0]], i32 0, i32 0
+// CHECK:         %[[VAL_21:.*]] = getelementptr inbounds [2 x ptr], ptr %[[VAL_1]], i32 0, i32 0
+// CHECK:         call void @__tgt_target_data_end_mapper(ptr @3, i64 -1, i32 2, ptr %[[VAL_20]], ptr %[[VAL_21]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
+// CHECK:         ret void
+
+// -----