[MLIR][OpenMP] Alterations to LLVM-IR lowering of OpenMP Dialect due to Bounds and MapEntry operations

[agozillon]

Summary

This patch adjusts the lower to LLVM-IR inside of
OpenMPToLLVMIRTranslation to faciliate the changes made
to Target related operations to add the new Map related
operations. It also includes adjustments to tests to support
these changes, primarily modifying the MLIR as opposed to
the LLVM-IR, the LLVM-IR should be identical after this patch.

Depends on D158735

[agozillon]

Patch 4/4 of the map changes from the patch series made by: https://reviews.llvm.org/D158732, https://reviews.llvm.org/D158735, https://reviews.llvm.org/D158737 and https://reviews.llvm.org/D158734 that aims to expand the current map support of the OpenMP dialect and lower from Fortran -> MLIR -> LLVM IR, future support on expanding the lowering from the OpenMP dialect to LLVM IR for map on TargetOp (omp target) will be forthcoming for declare target and explicit map variables hopefully in the near future (implicit likely a while after that).

This patch (provided I've setup the patch dependencies correctly) should pass CI, the others in the series will not.

[agozillon]

A small ping for some attention on this patch if a reviewer can spare some time, thank you very much!

[razvanlupusoru]

Looks reasonable to me based on the updated dialect design.

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

2210

This seems reasonable to me - but where is DataBoundsOp actually being handled? Or is that not implemented just yet?

[agozillon]

In D158737#4632382, @razvanlupusoru wrote:

Looks reasonable to me based on the updated dialect design.

Thank you very much for all your review points and input across the patch series, it's very appreciated.

Yes, this patch doesn't contain any altered lowering, it's just modifying the current status quo to support the dialect changes to map, the lowering will come in a later patch series provided this current series is successful. I thought 4 patches was already pushing it for all you poor reviewers (my apologies, and thank you all for the awesome work), as well as my own bandwidth!

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

2210

The handling of it is currently not implemented in this patch series. However, currently in a downstream patch (which I'll seek to upstream in pieces after/if this series lands) the information will be used alongside other information stored inside of MapEntryOp when lowering target operations that hold MapEntryOps (for the moment and near future it's only TargetOp, with the other operations on the horizon).

For example, in TargetOp where we perform map processing (now MapEntryOp processing), the lowering we have in place primarily uses the DataBoundsOp information to aid with array sectioning, so we primarily use the defined extents to generate some LLVM-IR operations differently (e.g. offset the array start and only pass so many elements).

Of course, if you have any pointers, ideas or directions I'd be more than happy to take them into consideration!

[agozillon]

rebase and update tests based on changes to the Map operations printing and parsing in seperate patch in patch series

[agozillon]

@razvanlupusoru thank you very much for your review!

I will leave the patch series open until Thursday to give time for further comments/feedback from other reviewers (in particular @kiranchandramohan and @jsjodin) at that point if no further comments/change requests are received I will push it upstream!

[TIFitis]

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

16

Can you please explain why the type is changing here?

[agozillon]

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

16

For offloading basic arguments the OpenMP runtime appears to also need the OMP_MAP_TARGET_PARAM bits set (more complex arguments need a variety of other bits set), at least for target just the tofrom/from/to bits aren't enough, perhaps it's different for target_data? I don't think that's the case, but I am perhaps wrong, I've only focused on target!

The target_param bits are being set in the OpenMPToLLVMTranslation phase at the moment as the maps are being processed, but I'm trying to move the setting of these bits to the frontend (the initial PFT lowering of map) where they can be kept together and reasoned about a bit better.

So it's changed to better represent what the IR lowering will be, with the addition of the target_param bits! I believe, currently in upstream only convertTargetOmp has the bits set during genMapInfos (and I believe I should actually include the removal of that segment in these patches actually, rather than future ones in hindsight, so I'll do that in the next update) for the moment.

[TIFitis]

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

16

As this affects correctness of generated code, I think we should be careful about making any changes that change the test output.

Can you please reason for why the map_type should be 35 here instead of 3 with examples from Clang perhaps? AFAIK the map_type should remain 3 here for this example. Similarly we might want to also look at the other tests with similar change.

[agozillon]

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

16

I was very incorrect! Thank you very much for pointing this out. I focused a little too much on the target case. It appears for target_data the map_types remain 3 (or the respective base mapping) and do not get appended with target_param, so I think will restrict the application of target_param to just TargetOp related map operations as it currently is (just later in the pipeline) and it can be expanded as we increase map support if required, if that sounds like a reasonable change?

For an example of the map_type value for Clang for target directives, this small program for target, when used in conjunction with the usual command, clang++ --offload-host-device --offload-arch=gfx90a -c -fopenmp -S -emit-llvm out-ir.cpp -o out-ir.ll :

#include <iostream>

int main() {
  int i = 1;

#pragma omp target map(tofrom : i)
{
    i = 99;
}

  std::cout << i << std::endl;

  return 0;
}

Should emit a map_type of:

@.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 35]

Currently, the change from 3 -> 35 (or 33/34 for to/from I believe) already occurs, but only for TargetOp's because of this change here: https://github.com/llvm/llvm-project/blob/main/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp#L1555

I'd like to move this to the frontend PFT lowering phase if possible, alongside where the rest of the map flags are currently set I believe, as I think it's likely better to keep them assigned in a similar location as best we can, so it's easier to reason about things (from the perspective of it mostly all being together and represented more clearly in the various stages of the IR). But perhaps I am wrong! :-)

It also could help to simplify this stage of the lowering a bit further, I think, for example, not all things need marked with OMP_TARGET_PARAM e.g. declare target variables don't require it, they require OMP_MAP_PTR_AND_OBJ instead, so we need to have a check that detects if something is declare target and then switch the appropriate bits. Implicit captures also need marked with OMP_IMPLICIT_MAP. Although, both of these checks can likely be done inside of the OpenMPToLLVMIRTranslation phase as well. I think it'd work perfectly fine in either location, at least from my current understanding. Although, at least some of the future added logic will still have to be inside of the OpenMPToLLVMIRTranslation phase as MEMBER_OF requires knowledge of the size of the combinedInfo arrays to set its bits (at least that appears to be how it's done in Clang).

Would you prefer us to continue setting the more specialized flags in the OpenMPToLLVMIRTranslation phase or does it sound reasonable to move what we can to the initial PFT lowering phase?

Thank you again for the great catch!

[TIFitis]

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

16

They way I see it, lowering phase is supposed to just translate the information as-is from high level code to FIR/MLIR.

A lot of the flags present here that can be set are specifically for the kernel launch and I don't think any MLIR passes etc benefit from having that information, rather it might just bloat the MLIR. In fact for cases such as implicitly captured operands, use_device_ptr/addr, privatisation etc. it might be infeasible to have all the flags correctly set at the lowering phase itself.

Hence my opinion would be to only lower the flags which are explicitly present in the code, and we can add the context based flags later on when lowering from MLIR to llvm.

Also, you are right. genMapInfos isn't complete at this point. The intention was to support the bare minimum of target data and target directives. As new directives are added or support for existing directives are expanded, we should also update genMapInfos accordingly to support the new cases.

[agozillon]

Update tests based on changes to the PFT lowering (no longer applies unneccesary flags, happens in later lowering in declare target patch and only for target as was happening before) and changes to the name of the operation and parsing/printing of the clauses.

[TIFitis]

LGTM :)

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

56

Nit: Can you please change the map type to from to prevent the change in test output here.

No need to push another review, feel free to commit with the change made.

[agozillon]

In D158737#4647321, @TIFitis wrote:

LGTM :)

Thank you for the reviews :-)

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

56

I will do so! I may have done an incorrect find and replace that caught this map clause at some point, thank you for the catch.

[agozillon]

@kiranchandramohan would it please be possible to get final acceptance on your end for each of the patches (https://reviews.llvm.org/D158732, https://reviews.llvm.org/D158734, https://reviews.llvm.org/D158735 and https://reviews.llvm.org/D158737) you've not signed off on if you're happy with them landing? (I have a minor buildbot error from what I believe is a missing change that I lost in a rebase that I plan on fixing today, but other than that no larger changes are forthcoming).

I'll land the patches after the additional acceptances provided no additional comments/change requests are made!

[kiranchandramohan]

In D158737#4647358, @agozillon wrote:

@kiranchandramohan would it please be possible to get final acceptance on your end for each of the patches (https://reviews.llvm.org/D158732, https://reviews.llvm.org/D158734, https://reviews.llvm.org/D158735 and https://reviews.llvm.org/D158737) you've not signed off on if you're happy with them landing? (I have a minor buildbot error from what I believe is a missing change that I lost in a rebase that I plan on fixing today, but other than that no larger changes are forthcoming).

I'll land the patches after the additional acceptances provided no additional comments/change requests are made!

OK. Please go ahead.

[agozillon]

• D158737 changes v3

[agozillon]

reverting the test back to its original form in most recent update as requested.

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

for (auto basePtr : combinedInfo.BasePointers) {
         return true;
       index++;
     }
     return false;
   };
 
   unsigned index = 0;
   for (const auto &mapOp : mapOperands) {
+    // Unlike dev_ptr and dev_addr operands these map operands point
+    // to a map entry operation which contains further information
+    // on the variable being mapped and how it should be mapped.
+    auto MapInfoOp =
+        mlir::dyn_cast<mlir::omp::MapInfoOp>(mapOp.getDefiningOp());
+
     // TODO: Only LLVMPointerTypes are handled.
-    if (!mapOp.getType().isa<LLVM::LLVMPointerType>())
+    if (!MapInfoOp.getType().isa<LLVM::LLVMPointerType>())
       return fail();
 
-    llvm::Value *mapOpValue = moduleTranslation.lookupValue(mapOp);
+    llvm::Value *mapOpValue =
+        moduleTranslation.lookupValue(MapInfoOp.getVarPtr());
     combinedInfo.BasePointers.emplace_back(mapOpValue);
     combinedInfo.Pointers.emplace_back(mapOpValue);
     combinedInfo.DevicePointers.emplace_back(
         llvm::OpenMPIRBuilder::DeviceInfoTy::None);
-    combinedInfo.Names.emplace_back(
-        LLVM::createMappingInformation(mapOp.getLoc(), *ompBuilder));
+    combinedInfo.Names.emplace_back(LLVM::createMappingInformation(
+        MapInfoOp.getVarPtr().getLoc(), *ompBuilder));
 
     combinedInfo.Types.emplace_back(
         llvm::omp::OpenMPOffloadMappingFlags(
-            mapTypes[index].dyn_cast<IntegerAttr>().getInt()) |
+            mapTypes[index].dyn_cast<IntegerAttr>().getUInt()) |
         (IsTargetParams
              ? llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM
              : llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_NONE));
 
     combinedInfo.Sizes.emplace_back(
-        builder.getInt64(getSizeInBytes(DL, mapOp.getType())));
+        builder.getInt64(getSizeInBytes(DL, MapInfoOp.getVarPtr().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();
convertOmpTargetData(Operation *op, llvm::IRBuilderBase &builder,
   SmallVector<Value> useDevPtrOperands;
   SmallVector<Value> useDevAddrOperands;
   ArrayAttr mapTypes;
   llvm::omp::RuntimeFunction RTLFn;
   DataLayout DL = DataLayout(op->getParentOfType<ModuleOp>());
 
   llvm::OpenMPIRBuilder *ompBuilder = moduleTranslation.getOpenMPBuilder();
 
+  auto getMapTypes = [](mlir::OperandRange mapOperands,
+                        mlir::MLIRContext *ctx) {
+    SmallVector<mlir::Attribute> mapTypes;
+    for (auto mapValue : mapOperands) {
+      if (mapValue.getDefiningOp()) {
+        auto mapOp =
+            mlir::dyn_cast<mlir::omp::MapInfoOp>(mapValue.getDefiningOp());
+        mapTypes.push_back(mapOp.getMapTypeAttr());
+      }
+    }
+    return mlir::ArrayAttr::get(ctx, mapTypes);
+  };
+
   LogicalResult result =
       llvm::TypeSwitch<Operation *, LogicalResult>(op)
           .Case([&](omp::DataOp dataOp) {
             if (auto ifExprVar = dataOp.getIfExpr())
               ifCond = moduleTranslation.lookupValue(ifExprVar);
 
             if (auto devId = dataOp.getDevice())
               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();
+            mapTypes =
+                getMapTypes(dataOp.getMapOperands(), dataOp->getContext());
             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 =
                       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();
+            mapTypes = getMapTypes(enterDataOp.getMapOperands(),
+                                   enterDataOp->getContext());
             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 =
                       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();
+            mapTypes = getMapTypes(exitDataOp.getMapOperands(),
+                                   exitDataOp->getContext());
             return success();
           })
           .Default([&](Operation *op) {
             return op->emitError("unsupported OpenMP operation: ")
                    << op->getName();
           });
 
   if (failed(result))
convertOmpTarget(Operation &opInst, llvm::IRBuilderBase &builder,
   int32_t defaultValTeams = -1;
   int32_t defaultValThreads = 0;
 
   llvm::OpenMPIRBuilder::InsertPointTy allocaIP =
       findAllocaInsertPoint(builder, moduleTranslation);
 
   DataLayout DL = DataLayout(opInst.getParentOfType<ModuleOp>());
   SmallVector<Value> mapOperands = targetOp.getMapOperands();
-  ArrayAttr mapTypes = targetOp.getMapTypes().value_or(nullptr);
+
+  auto getMapTypes = [](mlir::OperandRange mapOperands,
+                        mlir::MLIRContext *ctx) {
+    SmallVector<mlir::Attribute> mapTypes;
+    for (auto mapValue : mapOperands) {
+      if (mapValue.getDefiningOp()) {
+        auto mapOp =
+            mlir::dyn_cast<mlir::omp::MapInfoOp>(mapValue.getDefiningOp());
+        mapTypes.push_back(mapOp.getMapTypeAttr());
+      }
+    }
+    return mlir::ArrayAttr::get(ctx, mapTypes);
+  };
+
+  ArrayAttr mapTypes =
+      getMapTypes(targetOp.getMapOperands(), targetOp->getContext());
 
   llvm::OpenMPIRBuilder::MapInfosTy combinedInfos;
   auto genMapInfoCB = [&](llvm::OpenMPIRBuilder::InsertPointTy codeGenIP)
       -> llvm::OpenMPIRBuilder::MapInfosTy & {
     builder.restoreIP(codeGenIP);
     genMapInfos(builder, moduleTranslation, DL, combinedInfos, mapOperands,
                 mapTypes, {}, {}, true);
     return combinedInfos;
return llvm::TypeSwitch<Operation *, LogicalResult>(op)
         return convertOmpThreadprivate(*op, builder, moduleTranslation);
       })
       .Case<omp::DataOp, omp::EnterDataOp, omp::ExitDataOp>([&](auto op) {
         return convertOmpTargetData(op, builder, moduleTranslation);
       })
       .Case([&](omp::TargetOp) {
         return convertOmpTarget(*op, builder, moduleTranslation);
       })
+      .Case<omp::MapInfoOp, omp::DataBoundsOp>([&](auto op) {
+        // No-op, should be handled by relevant owning operations e.g.
+        // TargetOp, EnterDataOp, ExitDataOp, DataOp etc. and then
+        // discarded
+        return success();
+      })
       .Default([&](Operation *inst) {
         return inst->emitError("unsupported OpenMP operation: ")
                << inst->getName();
       });
 }
 
 void mlir::registerOpenMPDialectTranslation(DialectRegistry &registry) {
   registry.insert<omp::OpenMPDialect>();

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

 // RUN: mlir-translate -mlir-to-llvmir -split-input-file %s | FileCheck %s
 
 llvm.func @_QPopenmp_target_data() {
   %0 = llvm.mlir.constant(1 : i64) : i64
-  %1 = llvm.alloca %0 x i32 {bindc_name = "i", in_type = i32, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFopenmp_target_dataEi"} : (i64) -> !llvm.ptr<i32>
-  omp.target_data   map((tofrom -> %1 : !llvm.ptr<i32>)) {
-    %2 = llvm.mlir.constant(99 : i32) : i32
-    llvm.store %2, %1 : !llvm.ptr<i32>
+  %1 = llvm.alloca %0 x i32 {bindc_name = "i", in_type = i32, operand_segment_sizes = array<i32: 0, 0>, uniq_name = "_QFopenmp_target_dataEi"} : (i64) -> !llvm.ptr<i32>
+  %2 = omp.map_info var_ptr(%1 : !llvm.ptr<i32>)   map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
+  omp.target_data map_entries(%2 : !llvm.ptr<i32>) {
+    %3 = llvm.mlir.constant(99 : i32) : i32
+    llvm.store %3, %1 : !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 3]
 // CHECK-LABEL: define void @_QPopenmp_target_data() {
 // 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:         %[[VAL_11:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
 // CHECK:         %[[VAL_12:.*]] = 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_11]], ptr %[[VAL_12]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
 // CHECK:         ret void
 
 // -----
 
 llvm.func @_QPopenmp_target_data_region(%1 : !llvm.ptr<array<1024 x i32>>) {
-  omp.target_data   map((from -> %1 : !llvm.ptr<array<1024 x i32>>)) {
-    %2 = llvm.mlir.constant(99 : i32) : i32
-    %3 = llvm.mlir.constant(1 : i64) : i64
+  %2 = omp.map_info var_ptr(%1 : !llvm.ptr<array<1024 x i32>>)   map_clauses(from) capture(ByRef) -> !llvm.ptr<array<1024 x i32>> {name = ""}
+  omp.target_data map_entries(%2 : !llvm.ptr<array<1024 x i32>>) {
+    %3 = llvm.mlir.constant(99 : i32) : i32
     %4 = llvm.mlir.constant(1 : i64) : i64
-    %5 = llvm.mlir.constant(0 : i64) : i64
-    %6 = llvm.getelementptr %1[0, %5] : (!llvm.ptr<array<1024 x i32>>, i64) -> !llvm.ptr<i32>
-    llvm.store %2, %6 : !llvm.ptr<i32>
+    %5 = llvm.mlir.constant(1 : i64) : i64
+    %6 = llvm.mlir.constant(0 : i64) : i64
+    %7 = llvm.getelementptr %1[0, %6] : (!llvm.ptr<array<1024 x i32>>, i64) -> !llvm.ptr<i32>
+    llvm.store %3, %7 : !llvm.ptr<i32>
     omp.terminator
   }
   llvm.return
 }
 
 // CHECK:         @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 4096]
 // CHECK:         @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 2]
 // CHECK-LABEL: define void @_QPopenmp_target_data_region
 // CHECK:         (ptr %[[ARG_0:.*]]) {
 // 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:         br label %[[VAL_3:.*]]
 // CHECK:       entry:                                            ; preds = %[[VAL_4:.*]]
 // CHECK:         %[[VAL_5:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_0]], i32 0, i32 0
llvm.func @_QPomp_target_enter_exit(%1 : !llvm.ptr<array<1024 x i32>>, %3 : !llvm.ptr<array<512 x i32>>) {
   %8 = llvm.mlir.constant(5 : i32) : i32
   llvm.store %8, %7 : !llvm.ptr<i32>
   %9 = llvm.mlir.constant(2 : i32) : i32
   llvm.store %9, %5 : !llvm.ptr<i32>
   %10 = llvm.load %7 : !llvm.ptr<i32>
   %11 = llvm.mlir.constant(10 : i32) : i32
   %12 = llvm.icmp "slt" %10, %11 : i32
   %13 = llvm.load %5 : !llvm.ptr<i32>
-  omp.target_enter_data   if(%12 : i1) device(%13 : i32) map((to -> %1 : !llvm.ptr<array<1024 x i32>>), (alloc -> %3 : !llvm.ptr<array<512 x i32>>))
+  %map1 = omp.map_info var_ptr(%1 : !llvm.ptr<array<1024 x i32>>)   map_clauses(to) capture(ByRef) -> !llvm.ptr<array<1024 x i32>> {name = ""}
+  %map2 = omp.map_info var_ptr(%3 : !llvm.ptr<array<512 x i32>>)   map_clauses(exit_release_or_enter_alloc) capture(ByRef) -> !llvm.ptr<array<512 x i32>> {name = ""}
+  omp.target_enter_data   if(%12 : i1) device(%13 : i32) map_entries(%map1, %map2 : !llvm.ptr<array<1024 x i32>>, !llvm.ptr<array<512 x i32>>)
   %14 = llvm.load %7 : !llvm.ptr<i32>
   %15 = llvm.mlir.constant(10 : i32) : i32
   %16 = llvm.icmp "sgt" %14, %15 : i32
   %17 = llvm.load %5 : !llvm.ptr<i32>
-  omp.target_exit_data   if(%16 : i1) device(%17 : i32) map((from -> %1 : !llvm.ptr<array<1024 x i32>>), (release -> %3 : !llvm.ptr<array<512 x i32>>))
+  %map3 = omp.map_info var_ptr(%1 : !llvm.ptr<array<1024 x i32>>)   map_clauses(from) capture(ByRef) -> !llvm.ptr<array<1024 x i32>> {name = ""}
+  %map4 = omp.map_info var_ptr(%3 : !llvm.ptr<array<512 x i32>>)   map_clauses(exit_release_or_enter_alloc) capture(ByRef) -> !llvm.ptr<array<512 x i32>> {name = ""}
+  omp.target_exit_data   if(%16 : i1) device(%17 : i32) map_entries(%map3, %map4 : !llvm.ptr<array<1024 x i32>>, !llvm.ptr<array<512 x i32>>)
   llvm.return
 }
 
 // CHECK:         @.offload_sizes = private unnamed_addr constant [2 x i64] [i64 4096, i64 2048]
 // CHECK:         @.offload_maptypes = private unnamed_addr constant [2 x i64] [i64 1, i64 0]
 // CHECK:         @.offload_sizes.1 = private unnamed_addr constant [2 x i64] [i64 4096, i64 2048]
 // CHECK:         @.offload_maptypes.2 = private unnamed_addr constant [2 x i64] [i64 2, i64 0]
 // CHECK-LABEL: define void @_QPomp_target_enter_exit

 // 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>>)  {
+  %map1 = omp.map_info var_ptr(%a : !llvm.ptr<!llvm.ptr<i32>>)   map_clauses(from) capture(ByRef) -> !llvm.ptr<!llvm.ptr<i32>> {name = ""}
+  omp.target_data  map_entries(%map1 : !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:         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>>)  {
+  %map = omp.map_info var_ptr(%a : !llvm.ptr<!llvm.ptr<i32>>)   map_clauses(from) capture(ByRef) -> !llvm.ptr<!llvm.ptr<i32>> {name = ""}
+  omp.target_data  map_entries(%map : !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:         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>)  {
+  %map = omp.map_info var_ptr(%a : !llvm.ptr<i32>)   map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
+  omp.target_data  map_entries(%map : !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
 }
 

 
 // -----
 
 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>>)  {
+  %map = omp.map_info var_ptr(%b : !llvm.ptr<!llvm.ptr<i32>>)   map_clauses(from) capture(ByRef) -> !llvm.ptr<!llvm.ptr<i32>> {name = ""}
+  omp.target_data  map_entries(%map : !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.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>>)  {
+  %map = omp.map_info var_ptr(%a : !llvm.ptr<!llvm.ptr<i32>>)   map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<!llvm.ptr<i32>> {name = ""}
+  %map1 = omp.map_info var_ptr(%b : !llvm.ptr<!llvm.ptr<i32>>)   map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<!llvm.ptr<i32>> {name = ""}
+  omp.target_data  map_entries(%map, %map1 : !llvm.ptr<!llvm.ptr<i32>>, !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

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

 // RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
 
 module attributes {omp.is_target_device = true} {
   llvm.func @omp_target_region_() {
     %0 = llvm.mlir.constant(20 : i32) : i32
     %1 = llvm.mlir.constant(10 : i32) : i32
     %2 = llvm.mlir.constant(1 : i64) : i64
     %3 = llvm.alloca %2 x i32 {bindc_name = "a", in_type = i32, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEa"} : (i64) -> !llvm.ptr<i32>
     %4 = llvm.mlir.constant(1 : i64) : i64
     %5 = llvm.alloca %4 x i32 {bindc_name = "b", in_type = i32, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEb"} : (i64) -> !llvm.ptr<i32>
     %6 = llvm.mlir.constant(1 : i64) : i64
     %7 = llvm.alloca %6 x i32 {bindc_name = "c", in_type = i32, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEc"} : (i64) -> !llvm.ptr<i32>
     llvm.store %1, %3 : !llvm.ptr<i32>
     llvm.store %0, %5 : !llvm.ptr<i32>
-    omp.target map((to -> %3 : !llvm.ptr<i32>), (to -> %5 : !llvm.ptr<i32>), (from -> %7 : !llvm.ptr<i32>)) {
+    %map1 = omp.map_info var_ptr(%3 : !llvm.ptr<i32>)   map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
+    %map2 = omp.map_info var_ptr(%5 : !llvm.ptr<i32>)   map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
+    %map3 = omp.map_info var_ptr(%7 : !llvm.ptr<i32>)   map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
+    omp.target map_entries(%map1, %map2, %map3 : !llvm.ptr<i32>, !llvm.ptr<i32>, !llvm.ptr<i32>) {
       %8 = llvm.load %3 : !llvm.ptr<i32>
       %9 = llvm.load %5 : !llvm.ptr<i32>
       %10 = llvm.add %8, %9  : i32
       llvm.store %10, %7 : !llvm.ptr<i32>
       omp.terminator
     }
     llvm.return
   }

mlir/test/Target/LLVMIR/omptarget-region-llvm-target-device.mlir

 // This test checks that the name of the generated kernel function is using the
 // name stored in the omp.outline_parent_name attribute.
 // RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
 
 module attributes {omp.is_target_device = true} {
   llvm.func @writeindex_omp_outline_0_(%arg0: !llvm.ptr<i32>, %arg1: !llvm.ptr<i32>) attributes {omp.outline_parent_name = "writeindex_"} {
-    omp.target   map((from -> %arg0 : !llvm.ptr<i32>), (implicit -> %arg1: !llvm.ptr<i32>)) {
-      %0 = llvm.mlir.constant(20 : i32) : i32
-      %1 = llvm.mlir.constant(10 : i32) : i32
-      llvm.store %1, %arg0 : !llvm.ptr<i32>
-      llvm.store %0, %arg1 : !llvm.ptr<i32>
+    %0 = omp.map_info var_ptr(%arg0 : !llvm.ptr<i32>)   map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
+    %1 = omp.map_info var_ptr(%arg1 : !llvm.ptr<i32>)   map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
+    omp.target   map_entries(%0, %1 : !llvm.ptr<i32>, !llvm.ptr<i32>) {
+      %2 = llvm.mlir.constant(20 : i32) : i32
+      %3 = llvm.mlir.constant(10 : i32) : i32
+      llvm.store %3, %arg0 : !llvm.ptr<i32>
+      llvm.store %2, %arg1 : !llvm.ptr<i32>
       omp.terminator
     }
     llvm.return
   }
 }
 
-// CHECK: define {{.*}} void @__omp_offloading_{{.*}}_{{.*}}_writeindex__l7(ptr {{.*}}, ptr {{.*}}) {
+// CHECK: define {{.*}} void @__omp_offloading_{{.*}}_{{.*}}_writeindex__l{{.*}}(ptr {{.*}}, ptr {{.*}}) {

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

 // RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
 
 module attributes {omp.is_target_device = false} {
   llvm.func @omp_target_region_() {
     %0 = llvm.mlir.constant(20 : i32) : i32
     %1 = llvm.mlir.constant(10 : i32) : i32
     %2 = llvm.mlir.constant(1 : i64) : i64
     %3 = llvm.alloca %2 x i32 {bindc_name = "a", in_type = i32, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEa"} : (i64) -> !llvm.ptr<i32>
     %4 = llvm.mlir.constant(1 : i64) : i64
     %5 = llvm.alloca %4 x i32 {bindc_name = "b", in_type = i32, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEb"} : (i64) -> !llvm.ptr<i32>
     %6 = llvm.mlir.constant(1 : i64) : i64
     %7 = llvm.alloca %6 x i32 {bindc_name = "c", in_type = i32, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEc"} : (i64) -> !llvm.ptr<i32>
     llvm.store %1, %3 : !llvm.ptr<i32>
     llvm.store %0, %5 : !llvm.ptr<i32>
-    omp.target map((to -> %3 : !llvm.ptr<i32>), (to -> %5 : !llvm.ptr<i32>), (from -> %7 : !llvm.ptr<i32>)) {
+    %map1 = omp.map_info var_ptr(%3 : !llvm.ptr<i32>)   map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
+    %map2 = omp.map_info var_ptr(%5 : !llvm.ptr<i32>)   map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
+    %map3 = omp.map_info var_ptr(%7 : !llvm.ptr<i32>)   map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
+    omp.target map_entries(%map1, %map2, %map3 : !llvm.ptr<i32>, !llvm.ptr<i32>, !llvm.ptr<i32>) {
       %8 = llvm.load %3 : !llvm.ptr<i32>
       %9 = llvm.load %5 : !llvm.ptr<i32>
       %10 = llvm.add %8, %9  : i32
       llvm.store %10, %7 : !llvm.ptr<i32>
       omp.terminator
     }
     llvm.return
   }

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

 // RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
 
 module attributes {omp.is_target_device = false} {
   llvm.func @omp_target_region_() {
     %0 = llvm.mlir.constant(20 : i32) : i32
     %1 = llvm.mlir.constant(10 : i32) : i32
     %2 = llvm.mlir.constant(1 : i64) : i64
     %3 = llvm.alloca %2 x i32 {bindc_name = "a", in_type = i32, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEa"} : (i64) -> !llvm.ptr<i32>
     %4 = llvm.mlir.constant(1 : i64) : i64
     %5 = llvm.alloca %4 x i32 {bindc_name = "b", in_type = i32, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEb"} : (i64) -> !llvm.ptr<i32>
     %6 = llvm.mlir.constant(1 : i64) : i64
     %7 = llvm.alloca %6 x i32 {bindc_name = "c", in_type = i32, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFomp_target_regionEc"} : (i64) -> !llvm.ptr<i32>
     llvm.store %1, %3 : !llvm.ptr<i32>
     llvm.store %0, %5 : !llvm.ptr<i32>
-    omp.target map((to -> %3 : !llvm.ptr<i32>), (to -> %5 : !llvm.ptr<i32>), (from -> %7 : !llvm.ptr<i32>)) {
+    %map1 = omp.map_info var_ptr(%3 : !llvm.ptr<i32>)   map_clauses(tofrommap_info) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
+    %map2 = omp.map_info var_ptr(%5 : !llvm.ptr<i32>)   map_clauses(tofrommap_info) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
+    %map3 = omp.map_info var_ptr(%7 : !llvm.ptr<i32>)   map_clauses(tofrommap_info) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
+    omp.target map_entries( %map1, %map2, %map3 : !llvm.ptr<i32>, !llvm.ptr<i32>, !llvm.ptr<i32>) {
       omp.parallel {
         %8 = llvm.load %3 : !llvm.ptr<i32>
         %9 = llvm.load %5 : !llvm.ptr<i32>
         %10 = llvm.add %8, %9  : i32
         llvm.store %10, %7 : !llvm.ptr<i32>
         omp.terminator
         }
       omp.terminator