[flang-commits] [flang] [openmp] [mlir] [Flang][OpenMP] Initial mapping of Fortran pointers and allocatables for target devices (PR #71766)
via flang-commits
flang-commits at lists.llvm.org
Fri Jan 5 08:11:29 PST 2024
https://github.com/agozillon updated https://github.com/llvm/llvm-project/pull/71766
>From 30e917e8e96ffe662ae1f699d46c44e7185ae57d Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Fri, 1 Dec 2023 06:21:01 -0600
Subject: [PATCH 01/19] [Flang][OpenMP] Initial mapping of Fortran pointers and
allocatables for target devices
This patch seeks to add an initial lowering for pointers and allocatable variables explicitly captured by implicit and explicit map in Flang OpenMP.
---
.../include/flang/Optimizer/CodeGen/CodeGen.h | 7 +
flang/lib/Lower/OpenMP.cpp | 170 +++++++--
flang/lib/Optimizer/CodeGen/CMakeLists.txt | 1 +
flang/lib/Optimizer/CodeGen/CodeGen.cpp | 5 +
flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp | 87 +++++
.../OpenMP/map-types-and-sizes.f90 | 54 ++-
flang/test/Lower/OpenMP/FIR/array-bounds.f90 | 35 +-
flang/test/Lower/OpenMP/FIR/target.f90 | 3 +-
.../Lower/OpenMP/allocatable-array-bounds.f90 | 121 ++++++
flang/test/Lower/OpenMP/allocatable-map.f90 | 15 +
flang/test/Lower/OpenMP/array-bounds.f90 | 30 +-
flang/test/Lower/OpenMP/target.f90 | 3 +-
mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td | 16 +-
.../OpenMP/OpenMPToLLVMIRTranslation.cpp | 243 +++++++++++-
mlir/test/Dialect/OpenMP/ops.mlir | 14 +
...target-fortran-allocatable-types-host.mlir | 349 ++++++++++++++++++
.../allocatable-array-section-1d-bounds.f90 | 46 +++
.../allocatable-array-section-3d-bounds.f90 | 44 +++
.../fortran/allocatable-map-scopes.f95 | 66 ++++
.../failing/target_single_value_allocate.f90 | 27 ++
.../fortran/pointer-scopes-enter-exit-map.f90 | 83 +++++
...pointer-target-array-section-3d-bounds.f90 | 43 +++
.../fortran/pointer-target-map-scopes.f95 | 64 ++++
.../target_enter_exit_allocatables.f90 | 44 +++
.../fortran/target_enter_exit_array.f90 | 41 ++
25 files changed, 1519 insertions(+), 92 deletions(-)
create mode 100644 flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
create mode 100644 flang/test/Lower/OpenMP/allocatable-array-bounds.f90
create mode 100644 flang/test/Lower/OpenMP/allocatable-map.f90
create mode 100644 mlir/test/Target/LLVMIR/omptarget-fortran-allocatable-types-host.mlir
create mode 100644 openmp/libomptarget/test/offloading/fortran/allocatable-array-section-1d-bounds.f90
create mode 100644 openmp/libomptarget/test/offloading/fortran/allocatable-array-section-3d-bounds.f90
create mode 100644 openmp/libomptarget/test/offloading/fortran/allocatable-map-scopes.f95
create mode 100644 openmp/libomptarget/test/offloading/fortran/failing/target_single_value_allocate.f90
create mode 100644 openmp/libomptarget/test/offloading/fortran/pointer-scopes-enter-exit-map.f90
create mode 100644 openmp/libomptarget/test/offloading/fortran/pointer-target-array-section-3d-bounds.f90
create mode 100644 openmp/libomptarget/test/offloading/fortran/pointer-target-map-scopes.f95
create mode 100644 openmp/libomptarget/test/offloading/fortran/target_enter_exit_allocatables.f90
create mode 100644 openmp/libomptarget/test/offloading/fortran/target_enter_exit_array.f90
diff --git a/flang/include/flang/Optimizer/CodeGen/CodeGen.h b/flang/include/flang/Optimizer/CodeGen/CodeGen.h
index 5ea96c900bc630..27bea240fa8e8e 100644
--- a/flang/include/flang/Optimizer/CodeGen/CodeGen.h
+++ b/flang/include/flang/Optimizer/CodeGen/CodeGen.h
@@ -19,6 +19,7 @@
namespace fir {
struct NameUniquer;
+class LLVMTypeConverter;
#define GEN_PASS_DECL_FIRTOLLVMLOWERING
#define GEN_PASS_DECL_CODEGENREWRITE
@@ -80,6 +81,12 @@ std::unique_ptr<mlir::Pass> createLLVMDialectToLLVMPass(
std::unique_ptr<mlir::Pass> createBoxedProcedurePass();
std::unique_ptr<mlir::Pass> createBoxedProcedurePass(bool useThunks);
+/// Specialised conversion patterns of OpenMP operations for FIR to LLVM
+/// dialect, utilised in cases where the default OpenMP dialect handling cannot
+/// handle all cases for intermingled fir types and operations.
+void populateOpenMPFIRToLLVMConversionPatterns(
+ LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns);
+
// declarative passes
#define GEN_PASS_REGISTRATION
#include "flang/Optimizer/CodeGen/CGPasses.h.inc"
diff --git a/flang/lib/Lower/OpenMP.cpp b/flang/lib/Lower/OpenMP.cpp
index c3a570bf15ea0d..ae4e295f59625f 100644
--- a/flang/lib/Lower/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP.cpp
@@ -1718,30 +1718,82 @@ bool ClauseProcessor::processLink(
static mlir::omp::MapInfoOp
createMapInfoOp(fir::FirOpBuilder &builder, mlir::Location loc,
- mlir::Value baseAddr, std::stringstream &name,
- mlir::SmallVector<mlir::Value> bounds, uint64_t mapType,
- mlir::omp::VariableCaptureKind mapCaptureType,
- mlir::Type retTy) {
- mlir::Value varPtr, varPtrPtr;
- mlir::TypeAttr varType;
-
+ mlir::Value baseAddr, mlir::Value varPtrPtr, std::string name,
+ mlir::SmallVector<mlir::Value> bounds,
+ mlir::SmallVector<mlir::Value> members, uint64_t mapType,
+ mlir::omp::VariableCaptureKind mapCaptureType, mlir::Type retTy,
+ bool isVal = false) {
if (auto boxTy = baseAddr.getType().dyn_cast<fir::BaseBoxType>()) {
baseAddr = builder.create<fir::BoxAddrOp>(loc, baseAddr);
retTy = baseAddr.getType();
}
- varPtr = baseAddr;
- varType = mlir::TypeAttr::get(
+ mlir::TypeAttr varType = mlir::TypeAttr::get(
llvm::cast<mlir::omp::PointerLikeType>(retTy).getElementType());
mlir::omp::MapInfoOp op = builder.create<mlir::omp::MapInfoOp>(
- loc, retTy, varPtr, varType, varPtrPtr, bounds,
+ loc, retTy, baseAddr, varType, varPtrPtr, members, bounds,
builder.getIntegerAttr(builder.getIntegerType(64, false), mapType),
builder.getAttr<mlir::omp::VariableCaptureKindAttr>(mapCaptureType),
- builder.getStringAttr(name.str()));
+ builder.getStringAttr(name));
+
return op;
}
+static mlir::omp::MapInfoOp processDescriptorTypeMappings(
+ Fortran::semantics::SemanticsContext &semanticsContext,
+ Fortran::lower::StatementContext &stmtCtx,
+ Fortran::lower::AbstractConverter &converter, mlir::Location loc,
+ mlir::Value descriptorAddr, mlir::Value descDataBaseAddr,
+ mlir::SmallVector<mlir::Value> &bounds, std::string asFortran,
+ llvm::omp::OpenMPOffloadMappingFlags mapCaptureType) {
+ llvm::SmallVector<mlir::Value> descriptorBaseAddrMembers;
+ fir::FirOpBuilder &firOpBuilder = converter.getFirOpBuilder();
+
+ mlir::Value descriptor = descriptorAddr;
+ if (fir::isAssumedShape(fir::unwrapRefType(descriptorAddr.getType()))) {
+ mlir::OpBuilder::InsertPoint insPt = firOpBuilder.saveInsertionPoint();
+ firOpBuilder.setInsertionPointToStart(firOpBuilder.getAllocaBlock());
+ descriptor =
+ firOpBuilder.create<fir::AllocaOp>(loc, descriptorAddr.getType());
+ firOpBuilder.restoreInsertionPoint(insPt);
+ firOpBuilder.create<fir::StoreOp>(loc, descriptorAddr, descriptor);
+ }
+
+ mlir::Value baseAddrAddr = firOpBuilder.create<fir::BoxOffsetOp>(
+ loc, descriptor, fir::BoxFieldAttr::base_addr);
+
+ descriptorBaseAddrMembers.push_back(createMapInfoOp(
+ firOpBuilder, loc, descDataBaseAddr, baseAddrAddr, asFortran, bounds, {},
+ static_cast<std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
+ mapCaptureType),
+ mlir::omp::VariableCaptureKind::ByRef, descDataBaseAddr.getType()));
+
+ if (fir::BaseBoxType baseBoxTy = llvm::dyn_cast<fir::BaseBoxType>(
+ fir::unwrapRefType(descriptor.getType()))) {
+ if (fir::boxHasAddendum(baseBoxTy)) {
+ mlir::Value addendumAddrAddr = firOpBuilder.create<fir::BoxOffsetOp>(
+ loc, descriptor, fir::BoxFieldAttr::derived_type);
+ // NOTE: We have no addendumAddr, but we must provide a varPtr alongside
+ // the varPtrPtr so we replicate it across the two for now.
+ descriptorBaseAddrMembers.push_back(createMapInfoOp(
+ firOpBuilder, loc, addendumAddrAddr, addendumAddrAddr, asFortran, {},
+ {},
+ static_cast<
+ std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
+ mapCaptureType),
+ mlir::omp::VariableCaptureKind::ByRef, addendumAddrAddr.getType()));
+ }
+ }
+
+ return createMapInfoOp(
+ firOpBuilder, loc, descriptor, mlir::Value{}, asFortran, {},
+ descriptorBaseAddrMembers,
+ static_cast<std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
+ mapCaptureType),
+ mlir::omp::VariableCaptureKind::ByRef, descriptor.getType());
+}
+
bool ClauseProcessor::processMap(
mlir::Location currentLocation, const llvm::omp::Directive &directive,
Fortran::semantics::SemanticsContext &semanticsContext,
@@ -1801,28 +1853,41 @@ bool ClauseProcessor::processMap(
std::get<Fortran::parser::OmpObjectList>(mapClause->v.t).v) {
llvm::SmallVector<mlir::Value> bounds;
std::stringstream asFortran;
- Fortran::lower::AddrAndBoundsInfo info =
- Fortran::lower::gatherDataOperandAddrAndBounds<
- Fortran::parser::OmpObject, mlir::omp::DataBoundsOp,
- mlir::omp::DataBoundsType>(
- converter, firOpBuilder, semanticsContext, stmtCtx, ompObject,
- clauseLocation, asFortran, bounds, treatIndexAsSection);
- // Explicit map captures are captured ByRef by default,
- // optimisation passes may alter this to ByCopy or other capture
- // types to optimise
- mlir::Value mapOp = createMapInfoOp(
- firOpBuilder, clauseLocation, info.addr, asFortran, bounds,
- static_cast<
- std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
- mapTypeBits),
- mlir::omp::VariableCaptureKind::ByRef, info.addr.getType());
+ Fortran::lower::AddrAndBoundsInfo info = Fortran::lower::gatherDataOperandAddrAndBounds<
+ Fortran::parser::OmpObject, mlir::omp::DataBoundsType,
+ mlir::omp::DataBoundsOp>(
+ converter, firOpBuilder, semanticsContext, stmtCtx, ompObject,
+ clauseLocation, asFortran, bounds, treatIndexAsSection);
+ auto origSymbol =
+ converter.getSymbolAddress(*getOmpObjectSymbol(ompObject));
+ mlir::Value mapOp, symAddr;
+ if (fir::isPointerType(origSymbol.getType()) ||
+ fir::isAllocatableType(origSymbol.getType()) ||
+ fir::isAssumedShape(origSymbol.getType())) {
+ symAddr = origSymbol;
+ mapOp = processDescriptorTypeMappings(
+ semanticsContext, stmtCtx, converter, clauseLocation,
+ origSymbol, info.addr, bounds, asFortran.str(), mapTypeBits);
+ } else {
+ // Explicit map captures are captured ByRef by default,
+ // optimisation passes may alter this to ByCopy or other capture
+ // types to optimise
+ symAddr = info.addr;
+ mapOp = createMapInfoOp(
+ firOpBuilder, clauseLocation, info.addr, mlir::Value{},
+ asFortran.str(), bounds, {},
+ static_cast<std::underlying_type_t<
+ llvm::omp::OpenMPOffloadMappingFlags>>(mapTypeBits),
+ mlir::omp::VariableCaptureKind::ByRef, info.addr.getType());
+ }
mapOperands.push_back(mapOp);
if (mapSymTypes)
- mapSymTypes->push_back(info.addr.getType());
+ mapSymTypes->push_back(symAddr.getType());
if (mapSymLocs)
- mapSymLocs->push_back(info.addr.getLoc());
+ mapSymLocs->push_back(symAddr.getLoc());
+
if (mapSymbols)
mapSymbols->push_back(getOmpObjectSymbol(ompObject));
}
@@ -1928,12 +1993,24 @@ bool ClauseProcessor::processMotionClauses(
converter, firOpBuilder, semanticsContext, stmtCtx, ompObject,
clauseLocation, asFortran, bounds, treatIndexAsSection);
- mlir::Value mapOp = createMapInfoOp(
- firOpBuilder, clauseLocation, info.addr, asFortran, bounds,
- static_cast<
- std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
- mapTypeBits),
- mlir::omp::VariableCaptureKind::ByRef, info.addr.getType());
+ auto origSymbol =
+ converter.getSymbolAddress(*getOmpObjectSymbol(ompObject));
+ mlir::Value mapOp;
+ if (origSymbol && fir::isTypeWithDescriptor(origSymbol.getType())) {
+ mapOp = processDescriptorTypeMappings(
+ semanticsContext, stmtCtx, converter, clauseLocation,
+ origSymbol, info.addr, bounds, asFortran.str(), mapTypeBits);
+ } else {
+ // Explicit map captures are captured ByRef by default,
+ // optimisation passes may alter this to ByCopy or other capture
+ // types to optimise
+ mapOp = createMapInfoOp(
+ firOpBuilder, clauseLocation, info.addr, mlir::Value{},
+ asFortran.str(), bounds, {},
+ static_cast<std::underlying_type_t<
+ llvm::omp::OpenMPOffloadMappingFlags>>(mapTypeBits),
+ mlir::omp::VariableCaptureKind::ByRef, info.addr.getType());
+ }
mapOperands.push_back(mapOp);
}
@@ -2617,7 +2694,8 @@ static void genBodyOfTargetOp(
std::stringstream name;
firOpBuilder.setInsertionPoint(targetOp);
mlir::Value mapOp = createMapInfoOp(
- firOpBuilder, copyVal.getLoc(), copyVal, name, bounds,
+ firOpBuilder, copyVal.getLoc(), copyVal, mlir::Value{}, name.str(),
+ bounds, llvm::SmallVector<mlir::Value>{},
static_cast<
std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT),
@@ -2745,12 +2823,22 @@ genTargetOp(Fortran::lower::AbstractConverter &converter,
}
}
- mlir::Value mapOp = createMapInfoOp(
- converter.getFirOpBuilder(), baseOp.getLoc(), baseOp, name, bounds,
- static_cast<
- std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
- mapFlag),
- captureKind, baseOp.getType());
+ mlir::Value mapOp;
+ if (fir::isPointerType(baseOp.getType()) ||
+ fir::isAllocatableType(baseOp.getType()) ||
+ fir::isAssumedShape(baseOp.getType())) {
+ mapOp = processDescriptorTypeMappings(
+ semanticsContext, stmtCtx, converter, baseOp.getLoc(), baseOp,
+ baseAddr, bounds, name.str(), mapFlag);
+ } else {
+ mapOp = createMapInfoOp(
+ converter.getFirOpBuilder(), baseOp.getLoc(), baseOp,
+ mlir::Value{}, name.str(), bounds, {},
+ static_cast<
+ std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
+ mapFlag),
+ captureKind, baseOp.getType());
+ }
mapOperands.push_back(mapOp);
mapSymTypes.push_back(baseOp.getType());
diff --git a/flang/lib/Optimizer/CodeGen/CMakeLists.txt b/flang/lib/Optimizer/CodeGen/CMakeLists.txt
index 0daa97b00dfa00..b304b8c950b7f8 100644
--- a/flang/lib/Optimizer/CodeGen/CMakeLists.txt
+++ b/flang/lib/Optimizer/CodeGen/CMakeLists.txt
@@ -1,4 +1,5 @@
add_flang_library(FIRCodeGen
+ CodeGenOpenMP.cpp
BoxedProcedure.cpp
CGOps.cpp
CodeGen.cpp
diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index e07732d57880c5..9a45c88c828132 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -3882,6 +3882,11 @@ class FIRToLLVMLowering
mlir::populateMathToLibmConversionPatterns(pattern);
mlir::populateComplexToLLVMConversionPatterns(typeConverter, pattern);
mlir::populateVectorToLLVMConversionPatterns(typeConverter, pattern);
+
+ // Flang specific overloads for OpenMP operations, to allow for special
+ // handling of things like Box types.
+ fir::populateOpenMPFIRToLLVMConversionPatterns(typeConverter, pattern);
+
mlir::ConversionTarget target{*context};
target.addLegalDialect<mlir::LLVM::LLVMDialect>();
// The OpenMP dialect is legal for Operations without regions, for those
diff --git a/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp b/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
new file mode 100644
index 00000000000000..5de95ed807b3fb
--- /dev/null
+++ b/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
@@ -0,0 +1,87 @@
+#include "flang/Optimizer/CodeGen/CodeGen.h"
+
+#include "flang/Optimizer/Builder/FIRBuilder.h"
+#include "flang/Optimizer/Builder/LowLevelIntrinsics.h"
+#include "flang/Optimizer/Dialect/FIRDialect.h"
+#include "flang/Optimizer/Dialect/FIROps.h"
+#include "flang/Optimizer/Dialect/FIRType.h"
+#include "flang/Optimizer/Dialect/Support/FIRContext.h"
+#include "flang/Optimizer/Support/FatalError.h"
+#include "flang/Optimizer/Support/InternalNames.h"
+#include "mlir/Conversion/LLVMCommon/ConversionTarget.h"
+#include "mlir/Conversion/LLVMCommon/Pattern.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
+#include "mlir/IR/PatternMatch.h"
+#include "mlir/Transforms/DialectConversion.h"
+
+using namespace fir;
+
+#define DEBUG_TYPE "flang-codegen-openmp"
+
+// fir::LLVMTypeConverter for converting to LLVM IR dialect types.
+#include "flang/Optimizer/CodeGen/TypeConverter.h"
+
+namespace {
+/// A pattern that converts the region arguments in a single-region OpenMP
+/// operation to the LLVM dialect. The body of the region is not modified and is
+/// expected to either be processed by the conversion infrastructure or already
+/// contain ops compatible with LLVM dialect types.
+template <typename OpType>
+class OpenMPFIROpConversion : public mlir::ConvertOpToLLVMPattern<OpType> {
+public:
+ explicit OpenMPFIROpConversion(const fir::LLVMTypeConverter &lowering)
+ : mlir::ConvertOpToLLVMPattern<OpType>(lowering) {}
+
+ const fir::LLVMTypeConverter &lowerTy() const {
+ return *static_cast<const fir::LLVMTypeConverter *>(
+ this->getTypeConverter());
+ }
+};
+
+// FIR Op specific conversion for MapInfoOp that overwrites the default OpenMP
+// Dialect lowering, this allows FIR specific lowering of types, required for
+// descriptors of allocatables currently.
+struct MapInfoOpConversion
+ : public OpenMPFIROpConversion<mlir::omp::MapInfoOp> {
+ using OpenMPFIROpConversion::OpenMPFIROpConversion;
+
+ mlir::LogicalResult
+ matchAndRewrite(mlir::omp::MapInfoOp curOp, OpAdaptor adaptor,
+ mlir::ConversionPatternRewriter &rewriter) const override {
+ const mlir::TypeConverter *converter = getTypeConverter();
+ llvm::SmallVector<mlir::Type> resTypes;
+ if (failed(converter->convertTypes(curOp->getResultTypes(), resTypes)))
+ return mlir::failure();
+
+ llvm::SmallVector<mlir::NamedAttribute> newAttrs;
+ mlir::omp::MapInfoOp newOp;
+ for (mlir::NamedAttribute attr : curOp->getAttrs()) {
+ if (auto typeAttr = mlir::dyn_cast<mlir::TypeAttr>(attr.getValue())) {
+ mlir::Type newAttr;
+ if (fir::isPointerType(typeAttr.getValue()) ||
+ fir::isAllocatableType(typeAttr.getValue()) ||
+ fir::isAssumedShape(typeAttr.getValue())) {
+ newAttr = lowerTy().convertBoxTypeAsStruct(
+ mlir::cast<fir::BaseBoxType>(typeAttr.getValue()));
+ } else {
+ newAttr = converter->convertType(typeAttr.getValue());
+ }
+ newAttrs.emplace_back(attr.getName(), mlir::TypeAttr::get(newAttr));
+ } else {
+ newAttrs.push_back(attr);
+ }
+ }
+
+ rewriter.replaceOpWithNewOp<mlir::omp::MapInfoOp>(
+ curOp, resTypes, adaptor.getOperands(), newAttrs);
+
+ return mlir::success();
+ }
+};
+} // namespace
+
+void fir::populateOpenMPFIRToLLVMConversionPatterns(
+ LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns) {
+ patterns.add<MapInfoOpConversion>(converter);
+}
diff --git a/flang/test/Integration/OpenMP/map-types-and-sizes.f90 b/flang/test/Integration/OpenMP/map-types-and-sizes.f90
index f0a0e5e765b415..97723aae57d57c 100644
--- a/flang/test/Integration/OpenMP/map-types-and-sizes.f90
+++ b/flang/test/Integration/OpenMP/map-types-and-sizes.f90
@@ -30,7 +30,7 @@ subroutine mapType_array
!$omp end target
end subroutine mapType_array
-!CHECK: @.offload_sizes{{.*}} = private unnamed_addr constant [1 x i64] [i64 8]
+!CHECK: @.offload_sizes{{.*}} = private unnamed_addr constant [1 x i64] [i64 24]
!CHECK: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 547]
subroutine mapType_ptr
integer, pointer :: a
@@ -39,6 +39,37 @@ subroutine mapType_ptr
!$omp end target
end subroutine mapType_ptr
+!CHECK: @.offload_sizes{{.*}} = private unnamed_addr constant [1 x i64] [i64 24]
+!CHECK: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 547]
+subroutine mapType_allocatable
+ integer, allocatable :: a
+ allocate(a)
+ !$omp target
+ a = 10
+ !$omp end target
+ deallocate(a)
+end subroutine mapType_allocatable
+
+!CHECK: @.offload_sizes{{.*}} = private unnamed_addr constant [3 x i64] [i64 0, i64 24, i64 4]
+!CHECK: @.offload_maptypes{{.*}} = private unnamed_addr constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976710675]
+subroutine mapType_ptr_explicit
+ integer, pointer :: a
+ !$omp target map(tofrom: a)
+ a = 10
+ !$omp end target
+end subroutine mapType_ptr_explicit
+
+!CHECK: @.offload_sizes{{.*}} = private unnamed_addr constant [3 x i64] [i64 0, i64 24, i64 4]
+!CHECK: @.offload_maptypes{{.*}} = private unnamed_addr constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976710675]
+subroutine mapType_allocatable_explicit
+ integer, allocatable :: a
+ allocate(a)
+ !$omp target map(tofrom: a)
+ a = 10
+ !$omp end target
+ deallocate(a)
+end subroutine mapType_allocatable_explicit
+
!CHECK: @.offload_sizes{{.*}} = private unnamed_addr constant [2 x i64] [i64 8, i64 4]
!CHECK: @.offload_maptypes{{.*}} = private unnamed_addr constant [2 x i64] [i64 544, i64 800]
subroutine mapType_c_ptr
@@ -58,3 +89,24 @@ subroutine mapType_char
a = 'b'
!$omp end target
end subroutine mapType_char
+
+!CHECK-LABEL: define void @maptype_ptr_explicit_() {
+!CHECK: %[[ALLOCA:.*]] = alloca { ptr, i64, i32, i8, i8, i8, i8 }, i64 1, align 8
+!CHECK: %[[ALLOCA_GEP:.*]] = getelementptr { ptr, i64, i32, i8, i8, i8, i8 }, ptr %[[ALLOCA]], i32 1
+!CHECK: %[[ALLOCA_GEP_INT:.*]] = ptrtoint ptr %[[ALLOCA_GEP]] to i64
+!CHECK: %[[ALLOCA_INT:.*]] = ptrtoint ptr %[[ALLOCA]] to i64
+!CHECK: %[[SIZE_DIFF:.*]] = sub i64 %[[ALLOCA_GEP_INT]], %[[ALLOCA_INT]]
+!CHECK: %[[DIV:.*]] = sdiv exact i64 %[[SIZE_DIFF]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
+!CHECK: %[[OFFLOAD_SIZE_ARR:.*]] = getelementptr inbounds [3 x i64], ptr %.offload_sizes, i32 0, i32 0
+!CHECK: store i64 %[[DIV]], ptr %[[OFFLOAD_SIZE_ARR]], align 8
+
+
+!CHECK-LABEL: define void @maptype_allocatable_explicit_() {
+!CHECK: %[[ALLOCA:.*]] = alloca { ptr, i64, i32, i8, i8, i8, i8 }, i64 1, align 8
+!CHECK: %[[ALLOCA_GEP:.*]] = getelementptr { ptr, i64, i32, i8, i8, i8, i8 }, ptr %[[ALLOCA]], i32 1
+!CHECK: %[[ALLOCA_GEP_INT:.*]] = ptrtoint ptr %[[ALLOCA_GEP]] to i64
+!CHECK: %[[ALLOCA_INT:.*]] = ptrtoint ptr %[[ALLOCA]] to i64
+!CHECK: %[[SIZE_DIFF:.*]] = sub i64 %[[ALLOCA_GEP_INT]], %[[ALLOCA_INT]]
+!CHECK: %[[DIV:.*]] = sdiv exact i64 %[[SIZE_DIFF]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
+!CHECK: %[[OFFLOAD_SIZE_ARR:.*]] = getelementptr inbounds [3 x i64], ptr %.offload_sizes, i32 0, i32 0
+!CHECK: store i64 %[[DIV]], ptr %[[OFFLOAD_SIZE_ARR]], align 8
\ No newline at end of file
diff --git a/flang/test/Lower/OpenMP/FIR/array-bounds.f90 b/flang/test/Lower/OpenMP/FIR/array-bounds.f90
index 0e0aeaed0b553e..7b9802585f9425 100644
--- a/flang/test/Lower/OpenMP/FIR/array-bounds.f90
+++ b/flang/test/Lower/OpenMP/FIR/array-bounds.f90
@@ -35,6 +35,7 @@ module assumed_array_routines
contains
!ALL-LABEL: func.func @_QMassumed_array_routinesPassumed_shape_array(
!ALL-SAME: %[[ARG0:.*]]: !fir.box<!fir.array<?xi32>> {fir.bindc_name = "arr_read_write"})
+!ALL: %[[INTERMEDIATE_ALLOCA:.*]] = fir.alloca !fir.box<!fir.array<?xi32>>
!ALL: %[[ALLOCA:.*]] = fir.alloca i32 {bindc_name = "i", uniq_name = "_QMassumed_array_routinesFassumed_shape_arrayEi"}
!ALL: %[[C0:.*]] = arith.constant 1 : index
!ALL: %[[C1:.*]] = arith.constant 0 : index
@@ -44,23 +45,25 @@ module assumed_array_routines
!ALL: %[[C0_1:.*]] = arith.constant 0 : index
!ALL: %[[DIMS1:.*]]:3 = fir.box_dims %arg0, %[[C0_1]] : (!fir.box<!fir.array<?xi32>>, index) -> (index, index, index)
!ALL: %[[BOUNDS:.*]] = omp.bounds lower_bound(%[[C3]] : index) upper_bound(%[[C4]] : index) extent(%[[DIMS1]]#1 : index) stride(%[[DIMS0]]#2 : index) start_idx(%[[C0]] : index) {stride_in_bytes = true}
-!ALL: %[[ADDROF:.*]] = fir.box_addr %arg0 : (!fir.box<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
-!ALL: %[[MAP:.*]] = omp.map_info var_ptr(%[[ADDROF]] : !fir.ref<!fir.array<?xi32>>, !fir.array<?xi32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS]]) -> !fir.ref<!fir.array<?xi32>> {name = "arr_read_write(2:5)"}
+!ALL: %[[BOXADDRADDR:.*]] = fir.box_offset %0 base_addr : (!fir.ref<!fir.box<!fir.array<?xi32>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
+!ALL: %[[BOXADDR:.*]] = fir.box_addr %[[ARG0]] : (!fir.box<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
+!ALL: %[[MAP_MEMBER:.*]] = omp.map_info var_ptr(%[[BOXADDR]] : !fir.ref<!fir.array<?xi32>>, !fir.array<?xi32>) var_ptr_ptr(%[[BOXADDRADDR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS]]) -> !fir.ref<!fir.array<?xi32>> {name = "arr_read_write(2:5)"}
+!ALL: %[[MAP:.*]] = omp.map_info var_ptr(%0 : !fir.ref<!fir.box<!fir.array<?xi32>>>, !fir.box<!fir.array<?xi32>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_MEMBER]] : !fir.ref<!fir.array<?xi32>>) -> !fir.ref<!fir.box<!fir.array<?xi32>>> {name = "arr_read_write(2:5)"}
!ALL: %[[MAP2:.*]] = omp.map_info var_ptr(%[[ALLOCA]] : !fir.ref<i32>, i32) map_clauses(implicit, exit_release_or_enter_alloc) capture(ByCopy) -> !fir.ref<i32> {name = "i"}
-!ALL: omp.target map_entries(%[[MAP]] -> %{{.*}}, %[[MAP2]] -> %{{.*}} : !fir.ref<!fir.array<?xi32>>, !fir.ref<i32>) {
-
+!ALL: omp.target map_entries(%[[MAP]] -> %{{.*}}, %[[MAP2]] -> %{{.*}} : !fir.ref<!fir.box<!fir.array<?xi32>>>, !fir.ref<i32>) {
subroutine assumed_shape_array(arr_read_write)
- integer, intent(inout) :: arr_read_write(:)
+ integer, intent(inout) :: arr_read_write(:)
!$omp target map(tofrom:arr_read_write(2:5))
do i = 2, 5
arr_read_write(i) = i
end do
!$omp end target
- end subroutine assumed_shape_array
+ end subroutine assumed_shape_array
!ALL-LABEL: func.func @_QMassumed_array_routinesPassumed_size_array(
!ALL-SAME: %[[ARG0:.*]]: !fir.ref<!fir.array<?xi32>> {fir.bindc_name = "arr_read_write"})
+!ALL: %[[UNDEF:.*]] = fir.undefined index
!ALL: %[[ALLOCA:.*]] = fir.alloca i32 {bindc_name = "i", uniq_name = "_QMassumed_array_routinesFassumed_size_arrayEi"}
!ALL: %[[C0:.*]] = arith.constant 1 : index
!ALL: %[[C1:.*]] = arith.constant 1 : index
@@ -71,17 +74,16 @@ end subroutine assumed_shape_array
!ALL: %[[MAP:.*]] = omp.map_info var_ptr(%[[ARG0]] : !fir.ref<!fir.array<?xi32>>, !fir.array<?xi32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS]]) -> !fir.ref<!fir.array<?xi32>> {name = "arr_read_write(2:5)"}
!ALL: %[[MAP2:.*]] = omp.map_info var_ptr(%[[ALLOCA]] : !fir.ref<i32>, i32) map_clauses(implicit, exit_release_or_enter_alloc) capture(ByCopy) -> !fir.ref<i32> {name = "i"}
!ALL: omp.target map_entries(%[[MAP]] -> %{{.*}}, %[[MAP2]] -> %{{.*}} : !fir.ref<!fir.array<?xi32>>, !fir.ref<i32>) {
+ subroutine assumed_size_array(arr_read_write)
+ integer, intent(inout) :: arr_read_write(*)
- subroutine assumed_size_array(arr_read_write)
- integer, intent(inout) :: arr_read_write(*)
-
- !$omp target map(tofrom:arr_read_write(2:5))
- do i = 2, 5
- arr_read_write(i) = i
- end do
- !$omp end target
- end subroutine assumed_size_array
- end module assumed_array_routines
+ !$omp target map(tofrom:arr_read_write(2:5))
+ do i = 2, 5
+ arr_read_write(i) = i
+ end do
+ !$omp end target
+ end subroutine assumed_size_array
+end module assumed_array_routines
!DEVICE-NOT:func.func @_QPcall_assumed_shape_and_size_array() {
@@ -113,7 +115,6 @@ end module assumed_array_routines
!HOST:fir.call @_QMassumed_array_routinesPassumed_size_array(%[[ARG1]]) fastmath<contract> : (!fir.ref<!fir.array<?xi32>>) -> ()
!HOST:return
!HOST:}
-
subroutine call_assumed_shape_and_size_array
use assumed_array_routines
integer :: arr_read_write(20)
diff --git a/flang/test/Lower/OpenMP/FIR/target.f90 b/flang/test/Lower/OpenMP/FIR/target.f90
index 5d36699bf0e902..2018958b4e28cf 100644
--- a/flang/test/Lower/OpenMP/FIR/target.f90
+++ b/flang/test/Lower/OpenMP/FIR/target.f90
@@ -450,7 +450,8 @@ end subroutine omp_target_device_ptr
subroutine omp_target_device_addr
integer, pointer :: a
!CHECK: %[[VAL_0:.*]] = fir.alloca !fir.box<!fir.ptr<i32>> {bindc_name = "a", uniq_name = "_QFomp_target_device_addrEa"}
- !CHECK: %[[MAP:.*]] = omp.map_info var_ptr({{.*}}) map_clauses(tofrom) capture(ByRef) -> {{.*}} {name = "a"}
+ !CHECK: %[[MAP_MEMBERS:.*]] = omp.map_info var_ptr({{.*}} : !fir.ptr<i32>, i32) var_ptr_ptr({{.*}} : !fir.llvm_ptr<!fir.ref<i32>>) map_clauses(tofrom) capture(ByRef) -> !fir.ptr<i32> {name = "a"}
+ !CHECK: %[[MAP:.*]] = omp.map_info var_ptr({{.*}} : !fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.box<!fir.ptr<i32>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_MEMBERS]] : !fir.ptr<i32>) -> !fir.ref<!fir.box<!fir.ptr<i32>>> {name = "a"}
!CHECK: omp.target_data map_entries(%[[MAP]] : {{.*}}) use_device_addr(%[[VAL_0]] : !fir.ref<!fir.box<!fir.ptr<i32>>>) {
!$omp target data map(tofrom: a) use_device_addr(a)
!CHECK: ^bb0(%[[VAL_1:.*]]: !fir.ref<!fir.box<!fir.ptr<i32>>>):
diff --git a/flang/test/Lower/OpenMP/allocatable-array-bounds.f90 b/flang/test/Lower/OpenMP/allocatable-array-bounds.f90
new file mode 100644
index 00000000000000..5bf84dc3d5962a
--- /dev/null
+++ b/flang/test/Lower/OpenMP/allocatable-array-bounds.f90
@@ -0,0 +1,121 @@
+!RUN: %flang_fc1 -emit-hlfir -fopenmp %s -o - | FileCheck %s --check-prefixes HOST
+
+!HOST-LABEL: func.func @_QPread_write_section() {
+
+!HOST: %[[ALLOCA_1:.*]] = fir.alloca !fir.box<!fir.heap<!fir.array<?xi32>>> {bindc_name = "sp_read", uniq_name = "_QFread_write_sectionEsp_read"}
+!HOST: %[[DECLARE_1:.*]]:2 = hlfir.declare %[[ALLOCA_1]] {fortran_attrs = #fir.var_attrs<allocatable>, uniq_name = "_QFread_write_sectionEsp_read"} : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>)
+
+!HOST: %[[ALLOCA_2:.*]] = fir.alloca !fir.box<!fir.heap<!fir.array<?xi32>>> {bindc_name = "sp_write", uniq_name = "_QFread_write_sectionEsp_write"}
+!HOST: %[[DECLARE_2:.*]]:2 = hlfir.declare %[[ALLOCA_2]] {fortran_attrs = #fir.var_attrs<allocatable>, uniq_name = "_QFread_write_sectionEsp_write"} : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>)
+
+!HOST: %[[LOAD_1:.*]] = fir.load %[[DECLARE_1]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+!HOST: %[[LOAD_2:.*]] = fir.load %[[DECLARE_1]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+!HOST: %[[CONSTANT_1:.*]] = arith.constant 0 : index
+!HOST: %[[BOX_1:.*]]:3 = fir.box_dims %[[LOAD_2]], %[[CONSTANT_1]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>, index) -> (index, index, index)
+!HOST: %[[CONSTANT_2:.*]] = arith.constant 0 : index
+!HOST: %[[BOX_2:.*]]:3 = fir.box_dims %[[LOAD_1]], %[[CONSTANT_2]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>, index) -> (index, index, index)
+!HOST: %[[CONSTANT_3:.*]] = arith.constant 2 : index
+!HOST: %[[LB_1:.*]] = arith.subi %[[CONSTANT_3]], %[[BOX_1]]#0 : index
+!HOST: %[[CONSTANT_4:.*]] = arith.constant 5 : index
+!HOST: %[[UB_1:.*]] = arith.subi %[[CONSTANT_4]], %[[BOX_1]]#0 : index
+!HOST: %[[LOAD_3:.*]] = fir.load %[[DECLARE_1]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+!HOST: %[[CONSTANT_3:.*]] = arith.constant 0 : index
+!HOST: %[[BOX_3:.*]]:3 = fir.box_dims %[[LOAD_3]], %[[CONSTANT_3]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>, index) -> (index, index, index)
+!HOST: %[[BOUNDS_1:.*]] = omp.bounds lower_bound(%[[LB_1]] : index) upper_bound(%[[UB_1]] : index) extent(%[[BOX_3]]#1 : index) stride(%[[BOX_2]]#2 : index) start_idx(%[[BOX_1]]#0 : index) {stride_in_bytes = true}
+!HOST: %[[VAR_PTR_PTR:.*]] = fir.box_offset %[[DECLARE_1]]#1 base_addr : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
+!HOST: %[[VAR_PTR:.*]] = fir.box_addr %[[LOAD_1]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>) -> !fir.heap<!fir.array<?xi32>>
+!HOST: %[[MAP_INFO_MEMBER:.*]] = omp.map_info var_ptr(%[[VAR_PTR]] : !fir.heap<!fir.array<?xi32>>, !fir.array<?xi32>) var_ptr_ptr(%[[VAR_PTR_PTR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS_1]]) -> !fir.heap<!fir.array<?xi32>> {name = "sp_read(2:5)"}
+!HOST: %[[MAP_INFO_1:.*]] = omp.map_info var_ptr(%[[DECLARE_1]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, !fir.box<!fir.heap<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_INFO_MEMBER]] : !fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> {name = "sp_read(2:5)"}
+
+!HOST: %[[LOAD_3:.*]] = fir.load %[[DECLARE_2]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+!HOST: %[[LOAD_4:.*]] = fir.load %[[DECLARE_2]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+!HOST: %[[CONSTANT_5:.*]] = arith.constant 0 : index
+!HOST: %[[BOX_3:.*]]:3 = fir.box_dims %[[LOAD_4]], %[[CONSTANT_5]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>, index) -> (index, index, index)
+!HOST: %[[CONSTANT_6:.*]] = arith.constant 0 : index
+!HOST: %[[BOX_4:.*]]:3 = fir.box_dims %[[LOAD_3]], %[[CONSTANT_6]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>, index) -> (index, index, index)
+!HOST: %[[CONSTANT_7:.*]] = arith.constant 2 : index
+!HOST: %[[LB_2:.*]] = arith.subi %[[CONSTANT_7]], %[[BOX_3]]#0 : index
+!HOST: %[[CONSTANT_8:.*]] = arith.constant 5 : index
+!HOST: %[[UB_2:.*]] = arith.subi %[[CONSTANT_8]], %[[BOX_3]]#0 : index
+!HOST: %[[LOAD_5:.*]] = fir.load %[[DECLARE_2]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+!HOST: %[[CONSTANT_5:.*]] = arith.constant 0 : index
+!HOST: %[[BOX_5:.*]]:3 = fir.box_dims %[[LOAD_5]], %[[CONSTANT_5]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>, index) -> (index, index, index)
+!HOST: %[[BOUNDS_2:.*]] = omp.bounds lower_bound(%[[LB_2]] : index) upper_bound(%[[UB_2]] : index) extent(%[[BOX_5]]#1 : index) stride(%[[BOX_4]]#2 : index) start_idx(%[[BOX_3]]#0 : index) {stride_in_bytes = true}
+!HOST: %[[VAR_PTR_PTR:.*]] = fir.box_offset %[[DECLARE_2]]#1 base_addr : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
+!HOST: %[[VAR_PTR:.*]] = fir.box_addr %[[LOAD_3]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>) -> !fir.heap<!fir.array<?xi32>>
+!HOST: %[[MAP_INFO_MEMBER:.*]] = omp.map_info var_ptr(%[[VAR_PTR]] : !fir.heap<!fir.array<?xi32>>, !fir.array<?xi32>) var_ptr_ptr(%[[VAR_PTR_PTR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS_2]]) -> !fir.heap<!fir.array<?xi32>> {name = "sp_write(2:5)"}
+!HOST: %[[MAP_INFO_2:.*]] = omp.map_info var_ptr(%[[DECLARE_2]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, !fir.box<!fir.heap<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_INFO_MEMBER]] : !fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> {name = "sp_write(2:5)"}
+
+subroutine read_write_section()
+ integer, allocatable :: sp_read(:)
+ integer, allocatable :: sp_write(:)
+ allocate(sp_read(10))
+ allocate(sp_write(10))
+ sp_write = (/0,0,0,0,0,0,0,0,0,0/)
+ sp_read = (/1,2,3,4,5,6,7,8,9,10/)
+
+!$omp target map(tofrom:sp_read(2:5)) map(tofrom:sp_write(2:5))
+ do i = 2, 5
+ sp_write(i) = sp_read(i)
+ end do
+!$omp end target
+end subroutine read_write_section
+
+module assumed_allocatable_array_routines
+ contains
+
+!HOST-LABEL: func.func @_QMassumed_allocatable_array_routinesPassumed_shape_array(
+
+!HOST: %[[DECLARE:.*]]:2 = hlfir.declare %[[ARG:.*]] {fortran_attrs = #fir.var_attrs<allocatable, intent_inout>, uniq_name = "_QMassumed_allocatable_array_routinesFassumed_shape_arrayEarr_read_write"} : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>)
+!HOST: %[[LOAD_1:.*]] = fir.load %[[DECLARE]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+!HOST: %[[LOAD_2:.*]] = fir.load %[[DECLARE]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+!HOST: %[[CONSTANT_1:.*]] = arith.constant 0 : index
+!HOST: %[[BOX_1:.*]]:3 = fir.box_dims %[[LOAD_2]], %[[CONSTANT_1]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>, index) -> (index, index, index)
+!HOST: %[[CONSTANT_2:.*]] = arith.constant 0 : index
+!HOST: %[[BOX_2:.*]]:3 = fir.box_dims %[[LOAD_1]], %[[CONSTANT_2]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>, index) -> (index, index, index)
+!HOST: %[[CONSTANT_3:.*]] = arith.constant 2 : index
+!HOST: %[[LB:.*]] = arith.subi %[[CONSTANT_3]], %[[BOX_1]]#0 : index
+!HOST: %[[CONSTANT_4:.*]] = arith.constant 5 : index
+!HOST: %[[UB:.*]] = arith.subi %[[CONSTANT_4]], %[[BOX_1]]#0 : index
+!HOST: %[[LOAD_3:.*]] = fir.load %[[DECLARE]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+!HOST: %[[CONSTANT_3:.*]] = arith.constant 0 : index
+!HOST: %[[BOX_3:.*]]:3 = fir.box_dims %[[LOAD_3]], %[[CONSTANT_3]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>, index) -> (index, index, index)
+!HOST: %[[BOUNDS:.*]] = omp.bounds lower_bound(%[[LB]] : index) upper_bound(%[[UB]] : index) extent(%[[BOX_3]]#1 : index) stride(%[[BOX_2]]#2 : index) start_idx(%[[BOX_1]]#0 : index) {stride_in_bytes = true}
+!HOST: %[[VAR_PTR_PTR:.*]] = fir.box_offset %[[DECLARE]]#1 base_addr : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
+!HOST: %[[VAR_PTR:.*]] = fir.box_addr %[[LOAD_1]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>) -> !fir.heap<!fir.array<?xi32>>
+!HOST: %[[MAP_INFO_MEMBER:.*]] = omp.map_info var_ptr(%[[VAR_PTR]] : !fir.heap<!fir.array<?xi32>>, !fir.array<?xi32>) var_ptr_ptr(%[[VAR_PTR_PTR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS]]) -> !fir.heap<!fir.array<?xi32>> {name = "arr_read_write(2:5)"}
+!HOST: %[[MAP_INFO:.*]] = omp.map_info var_ptr(%[[DECLARE]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, !fir.box<!fir.heap<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_INFO_MEMBER]] : !fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> {name = "arr_read_write(2:5)"}
+
+ subroutine assumed_shape_array(arr_read_write)
+ integer, allocatable, intent(inout) :: arr_read_write(:)
+
+!$omp target map(tofrom:arr_read_write(2:5))
+ do i = 2, 5
+ arr_read_write(i) = i
+ end do
+!$omp end target
+end subroutine assumed_shape_array
+end module assumed_allocatable_array_routines
+
+!HOST-LABEL: func.func @_QPcall_assumed_shape_and_size_array() {
+!HOST: %[[ALLOCA:.*]] = fir.alloca !fir.box<!fir.heap<!fir.array<?xi32>>> {bindc_name = "arr_read_write", uniq_name = "_QFcall_assumed_shape_and_size_arrayEarr_read_write"}
+!HOST: %[[DECLARE:.*]]:2 = hlfir.declare %[[ALLOCA]] {fortran_attrs = #fir.var_attrs<allocatable>, uniq_name = "_QFcall_assumed_shape_and_size_arrayEarr_read_write"} : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>)
+!HOST: %[[ALLOCA_MEM:.*]] = fir.allocmem !fir.array<?xi32>, %{{.*}} {fir.must_be_heap = true, uniq_name = "_QFcall_assumed_shape_and_size_arrayEarr_read_write.alloc"}
+!HOST: %[[SHAPE:.*]] = fir.shape %{{.*}} : (index) -> !fir.shape<1>
+!HOST: %[[EMBOX:.*]] = fir.embox %[[ALLOCA_MEM]](%[[SHAPE]]) : (!fir.heap<!fir.array<?xi32>>, !fir.shape<1>) -> !fir.box<!fir.heap<!fir.array<?xi32>>>
+!HOST: fir.store %[[EMBOX]] to %[[DECLARE]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+!HOST: %[[LOAD:.*]] = fir.load %[[DECLARE]]#0 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+!HOST: %[[CONSTANT_1:.*]] = arith.constant 10 : index
+!HOST: %[[CONSTANT_2:.*]] = arith.constant 20 : index
+!HOST: %[[CONSTANT_3:.*]] = arith.constant 1 : index
+!HOST: %[[CONSTANT_4:.*]] = arith.constant 11 : index
+!HOST: %[[SHAPE:.*]] = fir.shape %[[CONSTANT_4]] : (index) -> !fir.shape<1>
+!HOST: %[[DESIGNATE:.*]] = hlfir.designate %[[LOAD]] (%[[CONSTANT_1]]:%[[CONSTANT_2]]:%[[CONSTANT_3]]) shape %[[SHAPE]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>, index, index, index, !fir.shape<1>) -> !fir.ref<!fir.array<11xi32>>
+!HOST: fir.call @_QPassumed_size_array(%[[DESIGNATE]]) fastmath<contract> : (!fir.ref<!fir.array<11xi32>>) -> ()
+subroutine call_assumed_shape_and_size_array
+ use assumed_allocatable_array_routines
+ integer, allocatable :: arr_read_write(:)
+ allocate(arr_read_write(20))
+ call assumed_size_array(arr_read_write(10:20))
+ deallocate(arr_read_write)
+end subroutine call_assumed_shape_and_size_array
\ No newline at end of file
diff --git a/flang/test/Lower/OpenMP/allocatable-map.f90 b/flang/test/Lower/OpenMP/allocatable-map.f90
new file mode 100644
index 00000000000000..5b8302459ca124
--- /dev/null
+++ b/flang/test/Lower/OpenMP/allocatable-map.f90
@@ -0,0 +1,15 @@
+!RUN: %flang_fc1 -emit-hlfir -fopenmp %s -o - | FileCheck %s --check-prefixes="HLFIRDIALECT"
+
+!HLFIRDIALECT: %[[POINTER:.*]]:2 = hlfir.declare %0 {fortran_attrs = #fir.var_attrs<pointer>, uniq_name = "_QFpointer_routineEpoint"} : (!fir.ref<!fir.box<!fir.ptr<i32>>>) -> (!fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.ref<!fir.box<!fir.ptr<i32>>>)
+!HLFIRDIALECT: %[[LOAD:.*]] = fir.load %[[POINTER]]#1 : !fir.ref<!fir.box<!fir.ptr<i32>>>
+!HLFIRDIALECT: %[[BOX_OFF:.*]] = fir.box_offset %[[POINTER]]#1 base_addr : (!fir.ref<!fir.box<!fir.ptr<i32>>>) -> !fir.llvm_ptr<!fir.ref<i32>>
+!HLFIRDIALECT: %[[BOX_ADDR:.*]] = fir.box_addr %[[LOAD]] : (!fir.box<!fir.ptr<i32>>) -> !fir.ptr<i32>
+!HLFIRDIALECT: %[[POINTER_MAP_MEMBER:.*]] = omp.map_info var_ptr(%[[BOX_ADDR]] : !fir.ptr<i32>, i32) var_ptr_ptr(%[[BOX_OFF]] : !fir.llvm_ptr<!fir.ref<i32>>) map_clauses(implicit, tofrom) capture(ByRef) -> !fir.ptr<i32> {name = "point"}
+!HLFIRDIALECT: %[[POINTER_MAP:.*]] = omp.map_info var_ptr(%[[POINTER]]#1 : !fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.box<!fir.ptr<i32>>) map_clauses(implicit, tofrom) capture(ByRef) members(%[[POINTER_MAP_MEMBER]] : !fir.ptr<i32>) -> !fir.ref<!fir.box<!fir.ptr<i32>>> {name = "point"}
+!HLFIRDIALECT: omp.target map_entries({{.*}}, %[[POINTER_MAP]] -> {{.*}} : {{.*}}, !fir.ref<!fir.box<!fir.ptr<i32>>>) {
+subroutine pointer_routine()
+ integer, pointer :: point
+!$omp target map(tofrom:pointer)
+ point = 1
+!$omp end target
+end subroutine pointer_routine
diff --git a/flang/test/Lower/OpenMP/array-bounds.f90 b/flang/test/Lower/OpenMP/array-bounds.f90
index 0fb5af188fa938..34d32468679b68 100644
--- a/flang/test/Lower/OpenMP/array-bounds.f90
+++ b/flang/test/Lower/OpenMP/array-bounds.f90
@@ -40,6 +40,7 @@ module assumed_array_routines
!HOST-LABEL: func.func @_QMassumed_array_routinesPassumed_shape_array(
!HOST-SAME: %[[ARG0:.*]]: !fir.box<!fir.array<?xi32>> {fir.bindc_name = "arr_read_write"}) {
+!HOST: %[[INTERMEDIATE_ALLOCA:.*]] = fir.alloca !fir.box<!fir.array<?xi32>>
!HOST: %[[ARG0_DECL:.*]]:2 = hlfir.declare %[[ARG0]] {fortran_attrs = #fir.var_attrs<intent_inout>, uniq_name = "_QMassumed_array_routinesFassumed_shape_arrayEarr_read_write"} : (!fir.box<!fir.array<?xi32>>) -> (!fir.box<!fir.array<?xi32>>, !fir.box<!fir.array<?xi32>>)
!HOST: %[[C0:.*]] = arith.constant 1 : index
!HOST: %[[C1:.*]] = arith.constant 0 : index
@@ -49,9 +50,11 @@ module assumed_array_routines
!HOST: %[[C0_1:.*]] = arith.constant 0 : index
!HOST: %[[DIMS1:.*]]:3 = fir.box_dims %[[ARG0_DECL]]#1, %[[C0_1]] : (!fir.box<!fir.array<?xi32>>, index) -> (index, index, index)
!HOST: %[[BOUNDS:.*]] = omp.bounds lower_bound(%[[C3]] : index) upper_bound(%[[C4]] : index) extent(%[[DIMS1]]#1 : index) stride(%[[DIMS0]]#2 : index) start_idx(%[[C0]] : index) {stride_in_bytes = true}
-!HOST: %[[ADDROF:.*]] = fir.box_addr %[[ARG0_DECL]]#1 : (!fir.box<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
-!HOST: %[[MAP:.*]] = omp.map_info var_ptr(%[[ADDROF]] : !fir.ref<!fir.array<?xi32>>, !fir.array<?xi32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS]]) -> !fir.ref<!fir.array<?xi32>> {name = "arr_read_write(2:5)"}
-!HOST: omp.target map_entries(%[[MAP]] -> %{{.*}}, {{.*}} -> {{.*}} : !fir.ref<!fir.array<?xi32>>, !fir.ref<i32>) {
+!HOST: %[[VAR_PTR_PTR:.*]] = fir.box_offset %0 base_addr : (!fir.ref<!fir.box<!fir.array<?xi32>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
+!HOST: %[[VAR_PTR:.*]] = fir.box_addr %[[ARG0_DECL]]#1 : (!fir.box<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
+!HOST: %[[MAP_INFO_MEMBER:.*]] = omp.map_info var_ptr(%[[VAR_PTR]] : !fir.ref<!fir.array<?xi32>>, !fir.array<?xi32>) var_ptr_ptr(%[[VAR_PTR_PTR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS]]) -> !fir.ref<!fir.array<?xi32>> {name = "arr_read_write(2:5)"}
+!HOST: %[[MAP:.*]] = omp.map_info var_ptr(%[[INTERMEDIATE_ALLOCA]] : !fir.ref<!fir.box<!fir.array<?xi32>>>, !fir.box<!fir.array<?xi32>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_INFO_MEMBER]] : !fir.ref<!fir.array<?xi32>>) -> !fir.ref<!fir.box<!fir.array<?xi32>>> {name = "arr_read_write(2:5)"}
+!HOST: omp.target map_entries(%[[MAP]] -> %{{.*}}, {{.*}} -> {{.*}} : !fir.ref<!fir.box<!fir.array<?xi32>>>, !fir.ref<i32>) {
subroutine assumed_shape_array(arr_read_write)
integer, intent(inout) :: arr_read_write(:)
@@ -60,7 +63,7 @@ subroutine assumed_shape_array(arr_read_write)
arr_read_write(i) = i
end do
!$omp end target
- end subroutine assumed_shape_array
+ end subroutine assumed_shape_array
!HOST-LABEL: func.func @_QMassumed_array_routinesPassumed_size_array(
@@ -76,17 +79,16 @@ end subroutine assumed_shape_array
!HOST: %[[BOUNDS:.*]] = omp.bounds lower_bound(%[[C1]] : index) upper_bound(%[[C2]] : index) extent(%[[EXT]] : index) stride(%[[C0]] : index) start_idx(%[[C0]] : index)
!HOST: %[[MAP:.*]] = omp.map_info var_ptr(%[[ARG0_DECL]]#1 : !fir.ref<!fir.array<?xi32>>, !fir.array<?xi32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS]]) -> !fir.ref<!fir.array<?xi32>> {name = "arr_read_write(2:5)"}
!HOST: omp.target map_entries(%[[MAP]] -> %{{.*}}, {{.*}} -> {{.*}} : !fir.ref<!fir.array<?xi32>>, !fir.ref<i32>) {
- subroutine assumed_size_array(arr_read_write)
- integer, intent(inout) :: arr_read_write(*)
-
- !$omp target map(tofrom:arr_read_write(2:5))
- do i = 2, 5
- arr_read_write(i) = i
- end do
- !$omp end target
- end subroutine assumed_size_array
- end module assumed_array_routines
+ subroutine assumed_size_array(arr_read_write)
+ integer, intent(inout) :: arr_read_write(*)
+ !$omp target map(tofrom:arr_read_write(2:5))
+ do i = 2, 5
+ arr_read_write(i) = i
+ end do
+ !$omp end target
+ end subroutine assumed_size_array
+end module assumed_array_routines
!HOST-LABEL:func.func @_QPcall_assumed_shape_and_size_array() {
!HOST: %[[C20:.*]] = arith.constant 20 : index
diff --git a/flang/test/Lower/OpenMP/target.f90 b/flang/test/Lower/OpenMP/target.f90
index 5ca3a08d8a8b20..51571b48645760 100644
--- a/flang/test/Lower/OpenMP/target.f90
+++ b/flang/test/Lower/OpenMP/target.f90
@@ -445,7 +445,8 @@ subroutine omp_target_device_addr
integer, pointer :: a
!CHECK: %[[VAL_0:.*]] = fir.alloca !fir.box<!fir.ptr<i32>> {bindc_name = "a", uniq_name = "_QFomp_target_device_addrEa"}
!CHECK: %[[VAL_0_DECL:.*]]:2 = hlfir.declare %0 {fortran_attrs = #fir.var_attrs<pointer>, uniq_name = "_QFomp_target_device_addrEa"} : (!fir.ref<!fir.box<!fir.ptr<i32>>>) -> (!fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.ref<!fir.box<!fir.ptr<i32>>>)
- !CHECK: %[[MAP:.*]] = omp.map_info var_ptr({{.*}}) map_clauses(tofrom) capture(ByRef) -> {{.*}} {name = "a"}
+ !CHECK: %[[MAP_MEMBERS:.*]] = omp.map_info var_ptr({{.*}} : !fir.ptr<i32>, i32) var_ptr_ptr({{.*}} : !fir.llvm_ptr<!fir.ref<i32>>) map_clauses(tofrom) capture(ByRef) -> !fir.ptr<i32> {name = "a"}
+ !CHECK: %[[MAP:.*]] = omp.map_info var_ptr({{.*}} : !fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.box<!fir.ptr<i32>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_MEMBERS]] : !fir.ptr<i32>) -> !fir.ref<!fir.box<!fir.ptr<i32>>> {name = "a"}
!CHECK: omp.target_data map_entries(%[[MAP]] : {{.*}}) use_device_addr(%[[VAL_0_DECL]]#1 : !fir.ref<!fir.box<!fir.ptr<i32>>>) {
!$omp target data map(tofrom: a) use_device_addr(a)
!CHECK: ^bb0(%[[VAL_1:.*]]: !fir.ref<!fir.box<!fir.ptr<i32>>>):
diff --git a/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td b/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
index b9989b335a2aef..89096650d84142 100644
--- a/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
+++ b/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
@@ -1148,6 +1148,7 @@ def MapInfoOp : OpenMP_Op<"map_info", [AttrSizedOperandSegments]> {
let arguments = (ins OpenMP_PointerLikeType:$var_ptr,
TypeAttr:$var_type,
Optional<OpenMP_PointerLikeType>:$var_ptr_ptr,
+ Variadic<OpenMP_PointerLikeType>:$members,
Variadic<DataBoundsType>:$bounds, /* rank-0 to rank-{n-1} */
OptionalAttr<UI64Attr>:$map_type,
OptionalAttr<VariableCaptureKindAttr>:$map_capture_type,
@@ -1187,13 +1188,17 @@ def MapInfoOp : OpenMP_Op<"map_info", [AttrSizedOperandSegments]> {
- `var_type`: The type of the variable to copy.
- `var_ptr_ptr`: Used when the variable copied is a member of a class, structure
or derived type and refers to the originating struct.
+ - `members`: Used to indicate mapped child members for the current MapInfoOp,
+ represented as other MapInfoOp's, utilised in cases where a parent structure
+ type and members of the structure type are being mapped at the same time.
+ For example: map(to: parent, parent->member, parent->member2[:10])
- `bounds`: Used when copying slices of array's, pointers or pointer members of
- objects (e.g. derived types or classes), indicates the bounds to be copied
- of the variable. When it's an array slice it is in rank order where rank 0
- is the inner-most dimension.
+ objects (e.g. derived types or classes), indicates the bounds to be copied
+ of the variable. When it's an array slice it is in rank order where rank 0
+ is the inner-most dimension.
- 'map_clauses': OpenMP map type for this map capture, for example: from, to and
- always. It's a bitfield composed of the OpenMP runtime flags stored in
- OpenMPOffloadMappingFlags.
+ always. It's a bitfield composed of the OpenMP runtime flags stored in
+ OpenMPOffloadMappingFlags.
- 'map_capture_type': Capture type for the variable e.g. this, byref, byvalue, byvla
this can affect how the variable is lowered.
- `name`: Holds the name of variable as specified in user clause (including bounds).
@@ -1205,6 +1210,7 @@ def MapInfoOp : OpenMP_Op<"map_info", [AttrSizedOperandSegments]> {
`var_ptr_ptr` `(` $var_ptr_ptr `:` type($var_ptr_ptr) `)`
| `map_clauses` `(` custom<MapClause>($map_type) `)`
| `capture` `(` custom<CaptureType>($map_capture_type) `)`
+ | `members` `(` $members `:` type($members) `)`
| `bounds` `(` $bounds `)`
) `->` type($omp_ptr) attr-dict
}];
diff --git a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
index 629584683f4991..15152cf6a8357a 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
@@ -1676,14 +1676,14 @@ uint64_t getArrayElementSizeInBits(LLVM::LLVMArrayType arrTy, DataLayout &dl) {
// This function is somewhat equivalent to Clang's getExprTypeSize inside of
// CGOpenMPRuntime.cpp.
llvm::Value *getSizeInBytes(DataLayout &dl, const mlir::Type &type,
- Operation *clauseOp, llvm::IRBuilderBase &builder,
+ Operation *clauseOp, llvm::Value *basePointer,
+ llvm::Type *baseType, llvm::IRBuilderBase &builder,
LLVM::ModuleTranslation &moduleTranslation) {
// utilising getTypeSizeInBits instead of getTypeSize as getTypeSize gives
// the size in inconsistent byte or bit format.
uint64_t underlyingTypeSzInBits = dl.getTypeSizeInBits(type);
- if (auto arrTy = llvm::dyn_cast_if_present<LLVM::LLVMArrayType>(type)) {
+ if (auto arrTy = llvm::dyn_cast_if_present<LLVM::LLVMArrayType>(type))
underlyingTypeSzInBits = getArrayElementSizeInBits(arrTy, dl);
- }
if (auto memberClause =
mlir::dyn_cast_if_present<mlir::omp::MapInfoOp>(clauseOp)) {
@@ -1728,13 +1728,13 @@ void collectMapDataFromMapOperands(MapInfoData &mapData,
LLVM::ModuleTranslation &moduleTranslation,
DataLayout &dl,
llvm::IRBuilderBase &builder) {
- for (mlir::Value mapValue : mapOperands) {
+ auto mapFill = [&](mlir::Value mapValue) {
assert(mlir::isa<mlir::omp::MapInfoOp>(mapValue.getDefiningOp()) &&
"missing map info operation or incorrect map info operation type");
if (auto mapOp = mlir::dyn_cast_if_present<mlir::omp::MapInfoOp>(
mapValue.getDefiningOp())) {
- mapData.OriginalValue.push_back(
- moduleTranslation.lookupValue(mapOp.getVarPtr()));
+ mapData.OriginalValue.push_back(moduleTranslation.lookupValue(
+ mapOp.getVarPtrPtr() ? mapOp.getVarPtrPtr() : mapOp.getVarPtr()));
mapData.Pointers.push_back(mapData.OriginalValue.back());
if (llvm::Value *refPtr =
@@ -1747,10 +1747,11 @@ void collectMapDataFromMapOperands(MapInfoData &mapData,
mapData.BasePointers.push_back(mapData.OriginalValue.back());
}
- mapData.Sizes.push_back(getSizeInBytes(dl, mapOp.getVarType(), mapOp,
- builder, moduleTranslation));
mapData.BaseType.push_back(
moduleTranslation.convertType(mapOp.getVarType()));
+ mapData.Sizes.push_back(getSizeInBytes(
+ dl, mapOp.getVarType(), mapOp, mapData.BasePointers.back(),
+ mapData.BaseType.back(), builder, moduleTranslation));
mapData.MapClause.push_back(mapOp.getOperation());
mapData.Types.push_back(
llvm::omp::OpenMPOffloadMappingFlags(mapOp.getMapType().value()));
@@ -1759,6 +1760,183 @@ void collectMapDataFromMapOperands(MapInfoData &mapData,
mapData.DevicePointers.push_back(
llvm::OpenMPIRBuilder::DeviceInfoTy::None);
}
+ };
+
+ // In the case of Fortran descriptors some members get added implicitly
+ // after the target region has been generated during CodeGen lowering
+ // which prevents them from being added trivially to the target region
+ // as map arguments, we must handle this case here by generating
+ // MapInfoData for them.
+ // TODO: Revisit this when implementing derived types explicit member
+ // mapping, we likely want to represent these identically to simplify
+ // the overall lowering
+ SmallVector<Value> mapMemberOperands;
+ for (size_t i = 0; i < mapOperands.size(); ++i) {
+ auto mapInfoOp =
+ mlir::dyn_cast<mlir::omp::MapInfoOp>(mapOperands[i].getDefiningOp());
+ for (auto members : mapInfoOp.getMembers()) {
+ if (!std::any_of(mapOperands.begin(), mapOperands.end(),
+ [&](auto mapOp) {
+ return mapOp.getDefiningOp() ==
+ members.getDefiningOp();
+ }) &&
+ !std::any_of(mapMemberOperands.begin(), mapMemberOperands.end(),
+ [&](auto mapOp) {
+ return mapOp.getDefiningOp() ==
+ members.getDefiningOp();
+ }))
+ mapMemberOperands.push_back(members);
+ }
+ }
+
+ for (mlir::Value mapValue : mapOperands)
+ mapFill(mapValue);
+
+ for (mlir::Value mapValue : mapMemberOperands)
+ mapFill(mapValue);
+}
+
+static void processMapWithMembersOf(
+ LLVM::ModuleTranslation &moduleTranslation, llvm::IRBuilderBase &builder,
+ llvm::OpenMPIRBuilder &ompBuilder, DataLayout &dl,
+ llvm::OpenMPIRBuilder::MapInfosTy &combinedInfo, MapInfoData &mapData,
+ uint64_t mapDataIndex, bool isTargetParams) {
+ auto parentClause =
+ mlir::dyn_cast<mlir::omp::MapInfoOp>(mapData.MapClause[mapDataIndex]);
+
+ ////////// First Parent Map Segment //////////
+ // Map the first segment of our structure
+ combinedInfo.Types.emplace_back(
+ isTargetParams
+ ? llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM
+ : llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_NONE);
+ combinedInfo.DevicePointers.emplace_back(
+ llvm::OpenMPIRBuilder::DeviceInfoTy::None);
+ combinedInfo.Names.emplace_back(LLVM::createMappingInformation(
+ mapData.MapClause[mapDataIndex]->getLoc(), ompBuilder));
+ combinedInfo.BasePointers.emplace_back(mapData.BasePointers[mapDataIndex]);
+ combinedInfo.Pointers.emplace_back(mapData.Pointers[mapDataIndex]);
+
+ // Calculate size of the parent object being mapped based on the
+ // addresses at runtime, highAddr - lowAddr = size. This of course
+ // doesn't factor in allocated data like pointers, hence the further
+ // processing of members specified by users, or in the case of
+ // Fortran pointers and allocatables, the mapping of the pointed to
+ // data by the descriptor (which itself, is a structure containing
+ // runtime information on the dynamically allocated data).
+ llvm::Value *lowAddr = builder.CreatePointerCast(
+ mapData.Pointers[mapDataIndex], builder.getPtrTy());
+ llvm::Value *highAddr = builder.CreatePointerCast(
+ builder.CreateConstGEP1_32(mapData.BaseType[mapDataIndex],
+ mapData.Pointers[mapDataIndex], 1),
+ builder.getPtrTy());
+ llvm::Value *size = builder.CreateIntCast(
+ builder.CreatePtrDiff(builder.getInt8Ty(), highAddr, lowAddr),
+ builder.getInt64Ty(),
+ /*isSigned=*/false);
+ combinedInfo.Sizes.push_back(size);
+
+ ////////// Second Parent Map Segment //////////
+ // This creates the initial MEMBER_OF mapping that consists of
+ // the parent/top level container (same as above effectively, except
+ // with a fixed initial compile time size and seperate maptype which
+ // indicates the true mape type (tofrom etc.) and that it is a part
+ // of a larger mapping and indicating the link between it and it's
+ // members that are also explicitly mapped).
+ llvm::omp::OpenMPOffloadMappingFlags mapFlag =
+ llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_TO;
+ if (isTargetParams)
+ mapFlag &= ~llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM;
+
+ llvm::omp::OpenMPOffloadMappingFlags memberOfFlag =
+ ompBuilder.getMemberOfFlag(combinedInfo.BasePointers.size() - 1);
+ ompBuilder.setCorrectMemberOfFlag(mapFlag, memberOfFlag);
+
+ combinedInfo.Types.emplace_back(mapFlag);
+ combinedInfo.DevicePointers.emplace_back(
+ llvm::OpenMPIRBuilder::DeviceInfoTy::None);
+ combinedInfo.Names.emplace_back(LLVM::createMappingInformation(
+ mapData.MapClause[mapDataIndex]->getLoc(), ompBuilder));
+ combinedInfo.BasePointers.emplace_back(mapData.BasePointers[mapDataIndex]);
+ combinedInfo.Pointers.emplace_back(mapData.Pointers[mapDataIndex]);
+ combinedInfo.Sizes.emplace_back(mapData.Sizes[mapDataIndex]);
+
+ ////////// Mapping of Members Segment //////////
+ for (auto mappedMembers : parentClause.getMembers()) {
+ auto memberClause =
+ mlir::dyn_cast<mlir::omp::MapInfoOp>(mappedMembers.getDefiningOp());
+ int memberDataIdx = -1;
+ for (size_t i = 0; i < mapData.MapClause.size(); ++i) {
+ if (mapData.MapClause[i] == memberClause)
+ memberDataIdx = i;
+ }
+
+ assert(memberDataIdx >= 0 && "could not find mapped member of structure");
+
+ // Same MemberOfFlag to indicate its link with parent and other members
+ // of, and we flag that it's part of a pointer and object coupling.
+ auto mapFlag =
+ llvm::omp::OpenMPOffloadMappingFlags(memberClause.getMapType().value());
+ mapFlag &= ~llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM;
+ ompBuilder.setCorrectMemberOfFlag(mapFlag, memberOfFlag);
+ mapFlag |= llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_PTR_AND_OBJ;
+ combinedInfo.Types.emplace_back(mapFlag);
+ combinedInfo.DevicePointers.emplace_back(
+ llvm::OpenMPIRBuilder::DeviceInfoTy::None);
+ combinedInfo.Names.emplace_back(
+ LLVM::createMappingInformation(memberClause.getLoc(), ompBuilder));
+
+ combinedInfo.BasePointers.emplace_back(mapData.BasePointers[memberDataIdx]);
+
+ std::vector<llvm::Value *> idx{builder.getInt64(0)};
+ llvm::Value *offsetAddress = nullptr;
+ if (!memberClause.getBounds().empty()) {
+ if (mapData.BaseType[memberDataIdx]->isArrayTy()) {
+ for (int i = memberClause.getBounds().size() - 1; i >= 0; --i) {
+ if (auto boundOp = mlir::dyn_cast_if_present<mlir::omp::DataBoundsOp>(
+ memberClause.getBounds()[i].getDefiningOp())) {
+ idx.push_back(
+ moduleTranslation.lookupValue(boundOp.getLowerBound()));
+ }
+ }
+ } else {
+ std::vector<llvm::Value *> dimensionIndexSizeOffset{
+ builder.getInt64(1)};
+ for (size_t i = 1; i < memberClause.getBounds().size(); ++i) {
+ if (auto boundOp = mlir::dyn_cast_if_present<mlir::omp::DataBoundsOp>(
+ memberClause.getBounds()[i].getDefiningOp())) {
+ dimensionIndexSizeOffset.push_back(builder.CreateMul(
+ moduleTranslation.lookupValue(boundOp.getExtent()),
+ dimensionIndexSizeOffset[i - 1]));
+ }
+ }
+
+ for (int i = memberClause.getBounds().size() - 1; i >= 0; --i) {
+ if (auto boundOp = mlir::dyn_cast_if_present<mlir::omp::DataBoundsOp>(
+ memberClause.getBounds()[i].getDefiningOp())) {
+ if (!offsetAddress)
+ offsetAddress = builder.CreateMul(
+ moduleTranslation.lookupValue(boundOp.getLowerBound()),
+ dimensionIndexSizeOffset[i]);
+ else
+ offsetAddress = builder.CreateAdd(
+ offsetAddress,
+ builder.CreateMul(
+ moduleTranslation.lookupValue(boundOp.getLowerBound()),
+ dimensionIndexSizeOffset[i]));
+ }
+ }
+ }
+ }
+
+ llvm::Value *memberIdx =
+ builder.CreateLoad(builder.getPtrTy(), mapData.Pointers[memberDataIdx]);
+ memberIdx = builder.CreateInBoundsGEP(
+ mapData.BaseType[memberDataIdx], memberIdx,
+ offsetAddress ? std::vector<llvm::Value *>{offsetAddress} : idx,
+ "member_idx");
+ combinedInfo.Pointers.emplace_back(memberIdx);
+ combinedInfo.Sizes.emplace_back(mapData.Sizes[memberDataIdx]);
}
}
@@ -1782,15 +1960,56 @@ static void genMapInfos(llvm::IRBuilderBase &builder,
combinedInfo.Names.clear();
};
+ llvm::SmallVector<size_t, 4> primaryMapIdx;
+ for (size_t i = 0; i < mapData.MapClause.size(); ++i) {
+ primaryMapIdx.push_back(i);
+ }
+
+ // TODO: Handle nested MembersOf, currently only cares about the first level
+ // of nesting (all that was relevant for Fortran descriptors), but a slight
+ // refactoring of mapInfoData to hold nestings or membersOf may be a better
+ // approach to simplify things.
+ for (size_t i = 0; i < mapData.MapClause.size(); ++i) {
+ auto mapInfoOp = mlir::dyn_cast<mlir::omp::MapInfoOp>(mapData.MapClause[i]);
+ for (auto member : mapInfoOp.getMembers()) {
+ for (size_t j = 0; j < primaryMapIdx.size(); j++) {
+ if (member.getDefiningOp() == mapData.MapClause[primaryMapIdx[j]]) {
+ primaryMapIdx.erase(&primaryMapIdx[j]);
+ j--;
+ }
+ }
+ }
+ }
+
// We operate under the assumption that all vectors that are
// required in MapInfoData are of equal lengths (either filled with
// default constructed data or appropiate information) so we can
// utilise the size from any component of MapInfoData, if we can't
// something is missing from the initial MapInfoData construction.
- for (size_t i = 0; i < mapData.MapClause.size(); ++i) {
+ for (unsigned long i : primaryMapIdx) {
+ auto mapInfoOp = mlir::dyn_cast<mlir::omp::MapInfoOp>(mapData.MapClause[i]);
+
+ bool isImplicit =
+ mapInfoOp.getMapType().value() &
+ static_cast<
+ std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
+ llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT);
+
+ // TODO: Look into more thoroughly handling implicit semantics, this
+ // mimics what Clang currently does with structures, which is map the
+ // structure and not the internals i.e. the descriptor without the pointer
+ // itself, so not a deep copy, which may be incorrect for
+ // allocatables/pointers, but allows explicit mapping and initial
+ // enter/exit handling
+ if (!isImplicit && !mapInfoOp.getMembers().empty()) {
+ processMapWithMembersOf(moduleTranslation, builder, *ompBuilder, dl,
+ combinedInfo, mapData, i, isTargetParams);
+ continue;
+ }
+
// Declare Target Mappings are excluded from being marked as
- // OMP_MAP_TARGET_PARAM as they are not passed as parameters, they're marked
- // with OMP_MAP_PTR_AND_OBJ instead.
+ // OMP_MAP_TARGET_PARAM as they are not passed as parameters, they're
+ // marked with OMP_MAP_PTR_AND_OBJ instead.
auto mapFlag = mapData.Types[i];
if (mapData.IsDeclareTarget[i])
mapFlag |= llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_PTR_AND_OBJ;
@@ -2424,7 +2643,7 @@ convertOmpTarget(Operation &opInst, llvm::IRBuilderBase &builder,
};
llvm::SmallVector<llvm::Value *, 4> kernelInput;
- for (size_t i = 0; i < mapData.MapClause.size(); ++i) {
+ for (size_t i = 0; i < mapOperands.size(); ++i) {
// declare target arguments are not passed to kernels as arguments
if (!mapData.IsDeclareTarget[i])
kernelInput.push_back(mapData.OriginalValue[i]);
diff --git a/mlir/test/Dialect/OpenMP/ops.mlir b/mlir/test/Dialect/OpenMP/ops.mlir
index 3d4f6435572f7f..ca92239b6fde96 100644
--- a/mlir/test/Dialect/OpenMP/ops.mlir
+++ b/mlir/test/Dialect/OpenMP/ops.mlir
@@ -2094,3 +2094,17 @@ func.func @omp_target_update_data (%if_cond : i1, %device : si32, %map1: memref<
return
}
+// CHECK-LABEL: omp_targets_is_allocatable
+// CHECK-SAME: (%[[ARG0:.*]]: !llvm.ptr, %[[ARG1:.*]]: !llvm.ptr)
+func.func @omp_targets_is_allocatable(%arg0: !llvm.ptr, %arg1: !llvm.ptr) -> () {
+ // CHECK: %[[MAP0:.*]] = omp.map_info var_ptr(%[[ARG0]] : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>) map_clauses(tofrom) capture(ByRef) -> !llvm.ptr {is_fortran_allocatable, name = ""}
+ %mapv1 = omp.map_info var_ptr(%arg0 : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>) map_clauses(tofrom) capture(ByRef) -> !llvm.ptr {is_fortran_allocatable, name = ""}
+ // CHECK: %[[MAP1:.*]] = omp.map_info var_ptr(%[[ARG1]] : !llvm.ptr, !llvm.array<10 x i32>) map_clauses(to) capture(ByCopy) -> !llvm.ptr {name = ""}
+ %mapv2 = omp.map_info var_ptr(%arg1 : !llvm.ptr, !llvm.array<10 x i32>) map_clauses(to) capture(ByCopy) -> !llvm.ptr {name = ""}
+ // CHECK: omp.target map_entries(%[[MAP0]] -> {{.*}}, %[[MAP1]] -> {{.*}} : !llvm.ptr, !llvm.ptr)
+ omp.target map_entries(%mapv1 -> %arg2, %mapv2 -> %arg3 : !llvm.ptr, !llvm.ptr) {
+ ^bb0(%arg2: !llvm.ptr, %arg3 : !llvm.ptr):
+ omp.terminator
+ }
+ return
+}
diff --git a/mlir/test/Target/LLVMIR/omptarget-fortran-allocatable-types-host.mlir b/mlir/test/Target/LLVMIR/omptarget-fortran-allocatable-types-host.mlir
new file mode 100644
index 00000000000000..dc3b230a7a55ee
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/omptarget-fortran-allocatable-types-host.mlir
@@ -0,0 +1,349 @@
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
+
+// This test checks the offload sizes, map types and base pointers and pointers
+// provided to the OpenMP kernel argument structure are correct when lowering
+// to LLVM-IR from MLIR when the fortran allocatables flag is switched on and
+// a fortran allocatable descriptor type is provided alongside the omp.map_info,
+// the test utilises mapping of array sections, full arrays and individual
+// allocated scalars.
+
+module attributes {omp.is_target_device = false} {
+ llvm.func @malloc(i64) -> !llvm.ptr
+ llvm.func @_QQmain() {
+ %0 = llvm.mlir.constant(1 : i32) : i32
+ %1 = llvm.alloca %0 x !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)> {alignment = 8 : i64} : (i32) -> !llvm.ptr
+ %2 = llvm.alloca %0 x !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)> {alignment = 8 : i64} : (i32) -> !llvm.ptr
+ %3 = llvm.alloca %0 x !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)> {alignment = 8 : i64} : (i32) -> !llvm.ptr
+ %4 = llvm.alloca %0 x !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)> {alignment = 8 : i64} : (i32) -> !llvm.ptr
+ %5 = llvm.alloca %0 x !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)> {alignment = 8 : i64} : (i32) -> !llvm.ptr
+ %6 = llvm.alloca %0 x !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)> {alignment = 8 : i64} : (i32) -> !llvm.ptr
+ %7 = llvm.mlir.constant(10 : index) : i64
+ %8 = llvm.mlir.constant(5 : index) : i64
+ %9 = llvm.mlir.constant(2 : index) : i64
+ %10 = llvm.mlir.constant(1 : index) : i64
+ %11 = llvm.mlir.constant(0 : index) : i64
+ %12 = llvm.mlir.addressof @_QFEfull_arr : !llvm.ptr
+ %13 = llvm.mlir.constant(1 : i64) : i64
+ %14 = llvm.alloca %13 x !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)> {bindc_name = "scalar"} : (i64) -> !llvm.ptr
+ %15 = llvm.mlir.zero : !llvm.ptr
+ %16 = llvm.mlir.constant(27 : i32) : i32
+ %17 = llvm.getelementptr %15[1] : (!llvm.ptr) -> !llvm.ptr, f32
+ %18 = llvm.ptrtoint %17 : !llvm.ptr to i64
+ %19 = llvm.mlir.undef : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %20 = llvm.insertvalue %18, %19[1] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %21 = llvm.mlir.constant(20180515 : i32) : i32
+ %22 = llvm.insertvalue %21, %20[2] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %23 = llvm.mlir.constant(0 : i32) : i32
+ %24 = llvm.trunc %23 : i32 to i8
+ %25 = llvm.insertvalue %24, %22[3] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %26 = llvm.trunc %16 : i32 to i8
+ %27 = llvm.insertvalue %26, %25[4] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %28 = llvm.mlir.constant(2 : i32) : i32
+ %29 = llvm.trunc %28 : i32 to i8
+ %30 = llvm.insertvalue %29, %27[5] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %31 = llvm.insertvalue %24, %30[6] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %32 = llvm.insertvalue %15, %31[0] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ llvm.store %32, %6 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>, !llvm.ptr
+ %33 = llvm.load %6 : !llvm.ptr -> !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ llvm.store %33, %14 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>, !llvm.ptr
+ %34 = llvm.mlir.addressof @_QFEsect_arr : !llvm.ptr
+ %35 = llvm.mul %18, %7 : i64
+ %36 = llvm.call @malloc(%35) {operandSegmentSizes = array<i32: 0, 1>, uniq_name = "_QFEfull_arr.alloc"} : (i64) -> !llvm.ptr
+ %37 = llvm.mlir.undef : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %38 = llvm.insertvalue %18, %37[1] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %39 = llvm.insertvalue %21, %38[2] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %40 = llvm.trunc %0 : i32 to i8
+ %41 = llvm.insertvalue %40, %39[3] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %42 = llvm.insertvalue %26, %41[4] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %43 = llvm.insertvalue %29, %42[5] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %44 = llvm.insertvalue %24, %43[6] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %45 = llvm.insertvalue %13, %44[7, 0, 0] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %46 = llvm.insertvalue %7, %45[7, 0, 1] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %47 = llvm.insertvalue %18, %46[7, 0, 2] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %48 = llvm.insertvalue %36, %47[0] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ llvm.store %48, %5 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>, !llvm.ptr
+ %49 = llvm.load %5 : !llvm.ptr -> !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ llvm.store %49, %12 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>, !llvm.ptr
+ %50 = llvm.getelementptr %15[1] : (!llvm.ptr) -> !llvm.ptr, i32
+ %51 = llvm.ptrtoint %50 : !llvm.ptr to i64
+ %52 = llvm.mul %51, %7 : i64
+ %53 = llvm.call @malloc(%52) {operandSegmentSizes = array<i32: 0, 1>, uniq_name = "_QFEsect_arr.alloc"} : (i64) -> !llvm.ptr
+ %54 = llvm.mlir.constant(9 : i32) : i32
+ %55 = llvm.insertvalue %51, %37[1] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %56 = llvm.insertvalue %21, %55[2] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %57 = llvm.insertvalue %40, %56[3] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %58 = llvm.trunc %54 : i32 to i8
+ %59 = llvm.insertvalue %58, %57[4] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %60 = llvm.insertvalue %29, %59[5] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %61 = llvm.insertvalue %24, %60[6] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %62 = llvm.insertvalue %13, %61[7, 0, 0] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %63 = llvm.insertvalue %7, %62[7, 0, 1] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %64 = llvm.insertvalue %51, %63[7, 0, 2] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %65 = llvm.insertvalue %53, %64[0] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ llvm.store %65, %4 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>, !llvm.ptr
+ %66 = llvm.load %4 : !llvm.ptr -> !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ llvm.store %66, %34 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>, !llvm.ptr
+ %67 = llvm.call @malloc(%18) {in_type = f32, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFEscalar.alloc"} : (i64) -> !llvm.ptr
+ %68 = llvm.insertvalue %67, %31[0] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ llvm.store %68, %3 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>, !llvm.ptr
+ %69 = llvm.load %3 : !llvm.ptr -> !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ llvm.store %69, %14 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>, !llvm.ptr
+ %70 = llvm.load %12 : !llvm.ptr -> !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ llvm.store %70, %2 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>, !llvm.ptr
+ %71 = llvm.getelementptr %2[0, 7, %11, 0] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %72 = llvm.load %71 : !llvm.ptr -> i64
+ %73 = llvm.getelementptr %2[0, 7, %11, 1] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %74 = llvm.load %73 : !llvm.ptr -> i64
+ %75 = llvm.getelementptr %2[0, 7, %11, 2] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %76 = llvm.load %75 : !llvm.ptr -> i64
+ %77 = llvm.sub %74, %10 : i64
+ %78 = omp.bounds lower_bound(%11 : i64) upper_bound(%77 : i64) extent(%74 : i64) stride(%76 : i64) start_idx(%72 : i64) {stride_in_bytes = true}
+ %79 = llvm.getelementptr %12[0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %80 = omp.map_info var_ptr(%79 : !llvm.ptr, f32) map_clauses(tofrom) capture(ByRef) bounds(%78) -> !llvm.ptr {name = "full_arr"}
+ %81 = omp.map_info var_ptr(%12 : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>) map_clauses(tofrom) capture(ByRef) members(%80 : !llvm.ptr) -> !llvm.ptr {name = "full_arr"}
+ %82 = llvm.load %34 : !llvm.ptr -> !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ llvm.store %82, %1 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>, !llvm.ptr
+ %83 = llvm.getelementptr %1[0, 7, %11, 0] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %84 = llvm.load %83 : !llvm.ptr -> i64
+ %85 = llvm.getelementptr %1[0, 7, %11, 1] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %86 = llvm.load %85 : !llvm.ptr -> i64
+ %87 = llvm.getelementptr %1[0, 7, %11, 2] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %88 = llvm.load %87 : !llvm.ptr -> i64
+ %89 = llvm.sub %9, %84 : i64
+ %90 = llvm.sub %8, %84 : i64
+ %91 = omp.bounds lower_bound(%89 : i64) upper_bound(%90 : i64) extent(%86 : i64) stride(%88 : i64) start_idx(%84 : i64) {stride_in_bytes = true}
+ %92 = llvm.getelementptr %34[0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %93 = omp.map_info var_ptr(%92 : !llvm.ptr, i32) map_clauses(tofrom) capture(ByRef) bounds(%91) -> !llvm.ptr {name = "sect_arr(2:5)"}
+ %94 = omp.map_info var_ptr(%34 : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>) map_clauses(tofrom) capture(ByRef) members(%93 : !llvm.ptr) -> !llvm.ptr {name = "sect_arr(2:5)"}
+ %95 = llvm.getelementptr %14[0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %96 = omp.map_info var_ptr(%95 : !llvm.ptr, f32) map_clauses(tofrom) capture(ByRef) -> !llvm.ptr {name = "scalar"}
+ %97 = omp.map_info var_ptr(%14 : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>) map_clauses(tofrom) capture(ByRef) members(%96 : !llvm.ptr) -> !llvm.ptr {name = "scalar"}
+ omp.target map_entries(%81 -> %arg0, %94 -> %arg1, %97 -> %arg2 : !llvm.ptr, !llvm.ptr, !llvm.ptr) {
+ ^bb0(%arg0: !llvm.ptr, %arg1: !llvm.ptr, %arg2: !llvm.ptr):
+ %98 = llvm.mlir.constant(1 : i32) : i32
+ %99 = llvm.alloca %98 x !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)> {alignment = 8 : i64} : (i32) -> !llvm.ptr
+ %100 = llvm.alloca %98 x !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)> {alignment = 8 : i64} : (i32) -> !llvm.ptr
+ %101 = llvm.alloca %98 x !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)> {alignment = 8 : i64} : (i32) -> !llvm.ptr
+ %102 = llvm.mlir.constant(13 : i32) : i32
+ %103 = llvm.mlir.constant(0 : index) : i64
+ %104 = llvm.mlir.constant(3 : index) : i64
+ %105 = llvm.mlir.constant(1 : index) : i64
+ %106 = llvm.mlir.constant(1.000000e+00 : f32) : f32
+ %107 = llvm.load %arg0 : !llvm.ptr -> !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ llvm.store %107, %101 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>, !llvm.ptr
+ %108 = llvm.getelementptr %101[0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %109 = llvm.load %108 : !llvm.ptr -> !llvm.ptr
+ %110 = llvm.getelementptr %101[0, 7, %103, 0] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %111 = llvm.load %110 : !llvm.ptr -> i64
+ %112 = llvm.getelementptr %101[0, 7, %103, 1] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %113 = llvm.load %112 : !llvm.ptr -> i64
+ %114 = llvm.getelementptr %101[0, 7, %103, 2] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %115 = llvm.mlir.constant(1 : i64) : i64
+ %116 = llvm.mlir.constant(0 : i64) : i64
+ %117 = llvm.sub %105, %111 : i64
+ %118 = llvm.mul %117, %115 : i64
+ %119 = llvm.mul %118, %115 : i64
+ %120 = llvm.add %119, %116 : i64
+ %121 = llvm.getelementptr %109[%120] : (!llvm.ptr, i64) -> !llvm.ptr, f32
+ llvm.store %106, %121 : f32, !llvm.ptr
+ %122 = llvm.load %arg1 : !llvm.ptr -> !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ llvm.store %122, %100 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>, !llvm.ptr
+ %123 = llvm.getelementptr %100[0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %124 = llvm.load %123 : !llvm.ptr -> !llvm.ptr
+ %125 = llvm.getelementptr %100[0, 7, %103, 0] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %126 = llvm.load %125 : !llvm.ptr -> i64
+ %127 = llvm.getelementptr %100[0, 7, %103, 1] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %128 = llvm.load %127 : !llvm.ptr -> i64
+ %129 = llvm.getelementptr %100[0, 7, %103, 2] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %130 = llvm.sub %104, %126 : i64
+ %131 = llvm.mul %130, %115 : i64
+ %132 = llvm.mul %131, %115 : i64
+ %133 = llvm.add %132, %116 : i64
+ %134 = llvm.getelementptr %124[%133] : (!llvm.ptr, i64) -> !llvm.ptr, i32
+ llvm.store %98, %134 : i32, !llvm.ptr
+ %135 = llvm.sitofp %98 : i32 to f32
+ %136 = llvm.alloca %115 x f32 : (i64) -> !llvm.ptr
+ llvm.store %135, %136 : f32, !llvm.ptr
+ %137 = llvm.mlir.constant(27 : i32) : i32
+ %138 = llvm.mlir.zero : !llvm.ptr
+ %139 = llvm.getelementptr %138[1] : (!llvm.ptr) -> !llvm.ptr, f32
+ %140 = llvm.ptrtoint %139 : !llvm.ptr to i64
+ %141 = llvm.mlir.undef : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %142 = llvm.insertvalue %140, %141[1] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %143 = llvm.mlir.constant(20180515 : i32) : i32
+ %144 = llvm.insertvalue %143, %142[2] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %145 = llvm.mlir.constant(0 : i32) : i32
+ %146 = llvm.trunc %145 : i32 to i8
+ %147 = llvm.insertvalue %146, %144[3] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %148 = llvm.trunc %137 : i32 to i8
+ %149 = llvm.insertvalue %148, %147[4] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %150 = llvm.insertvalue %146, %149[5] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %151 = llvm.insertvalue %146, %150[6] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ %152 = llvm.insertvalue %136, %151[0] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
+ llvm.store %152, %99 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>, !llvm.ptr
+ %153 = llvm.mlir.addressof @_QQclX6606f061cfd55e600f1165aa97a647a5 : !llvm.ptr
+ %154 = llvm.call @_FortranAAssign(%arg2, %99, %153, %102) : (!llvm.ptr, !llvm.ptr, !llvm.ptr, i32) -> !llvm.struct<()>
+ omp.terminator
+ }
+ llvm.return
+ }
+ llvm.mlir.global internal @_QFEfull_arr() {addr_space = 0 : i32} : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)> {
+ %0 = llvm.mlir.constant(0 : index) : i64
+ %1 = llvm.mlir.zero : !llvm.ptr
+ %2 = llvm.mlir.constant(27 : i32) : i32
+ %3 = llvm.getelementptr %1[1] : (!llvm.ptr) -> !llvm.ptr, f32
+ %4 = llvm.ptrtoint %3 : !llvm.ptr to i64
+ %5 = llvm.mlir.undef : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %6 = llvm.insertvalue %4, %5[1] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %7 = llvm.mlir.constant(20180515 : i32) : i32
+ %8 = llvm.insertvalue %7, %6[2] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %9 = llvm.mlir.constant(1 : i32) : i32
+ %10 = llvm.trunc %9 : i32 to i8
+ %11 = llvm.insertvalue %10, %8[3] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %12 = llvm.trunc %2 : i32 to i8
+ %13 = llvm.insertvalue %12, %11[4] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %14 = llvm.mlir.constant(2 : i32) : i32
+ %15 = llvm.trunc %14 : i32 to i8
+ %16 = llvm.insertvalue %15, %13[5] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %17 = llvm.mlir.constant(0 : i32) : i32
+ %18 = llvm.trunc %17 : i32 to i8
+ %19 = llvm.insertvalue %18, %16[6] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %20 = llvm.mlir.constant(1 : i64) : i64
+ %21 = llvm.insertvalue %20, %19[7, 0, 0] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %22 = llvm.insertvalue %0, %21[7, 0, 1] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %23 = llvm.insertvalue %4, %22[7, 0, 2] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %24 = llvm.insertvalue %1, %23[0] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ llvm.return %24 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ }
+ llvm.mlir.global internal @_QFEsect_arr() {addr_space = 0 : i32} : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)> {
+ %0 = llvm.mlir.constant(0 : index) : i64
+ %1 = llvm.mlir.zero : !llvm.ptr
+ %2 = llvm.mlir.constant(9 : i32) : i32
+ %3 = llvm.getelementptr %1[1] : (!llvm.ptr) -> !llvm.ptr, i32
+ %4 = llvm.ptrtoint %3 : !llvm.ptr to i64
+ %5 = llvm.mlir.undef : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %6 = llvm.insertvalue %4, %5[1] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %7 = llvm.mlir.constant(20180515 : i32) : i32
+ %8 = llvm.insertvalue %7, %6[2] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %9 = llvm.mlir.constant(1 : i32) : i32
+ %10 = llvm.trunc %9 : i32 to i8
+ %11 = llvm.insertvalue %10, %8[3] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %12 = llvm.trunc %2 : i32 to i8
+ %13 = llvm.insertvalue %12, %11[4] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %14 = llvm.mlir.constant(2 : i32) : i32
+ %15 = llvm.trunc %14 : i32 to i8
+ %16 = llvm.insertvalue %15, %13[5] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %17 = llvm.mlir.constant(0 : i32) : i32
+ %18 = llvm.trunc %17 : i32 to i8
+ %19 = llvm.insertvalue %18, %16[6] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %20 = llvm.mlir.constant(1 : i64) : i64
+ %21 = llvm.insertvalue %20, %19[7, 0, 0] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %22 = llvm.insertvalue %0, %21[7, 0, 1] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %23 = llvm.insertvalue %4, %22[7, 0, 2] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %24 = llvm.insertvalue %1, %23[0] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ llvm.return %24 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ }
+ llvm.mlir.global linkonce constant @_QQclX6606f061cfd55e600f1165aa97a647a5() comdat(@__llvm_comdat::@_QQclX6606f061cfd55e600f1165aa97a647a5) {addr_space = 0 : i32} : !llvm.array<97 x i8> {
+ %0 = llvm.mlir.constant("/home/agozillo/git/flang-dev/work-dir/declare-target-map/omptarget-fortran-allocatables-test.f90\00") : !llvm.array<97 x i8>
+ llvm.return %0 : !llvm.array<97 x i8>
+ }
+ llvm.comdat @__llvm_comdat {
+ llvm.comdat_selector @_QQclX6606f061cfd55e600f1165aa97a647a5 any
+ }
+ llvm.mlir.global external constant @_QQEnvironmentDefaults() {addr_space = 0 : i32} : !llvm.ptr {
+ %0 = llvm.mlir.zero : !llvm.ptr
+ llvm.return %0 : !llvm.ptr
+ }
+ llvm.func @_FortranAAssign(!llvm.ptr, !llvm.ptr, !llvm.ptr, i32) -> !llvm.struct<()> attributes {sym_visibility = "private"}
+}
+
+// CHECK: @.offload_sizes = private unnamed_addr constant [9 x i64] [i64 0, i64 48, i64 0, i64 0, i64 48, i64 0, i64 0, i64 24, i64 4]
+// CHECK: @.offload_maptypes = private unnamed_addr constant [9 x i64] [i64 32, i64 281474976710657, i64 281474976710675, i64 32, i64 1125899906842625, i64 1125899906842643, i64 32, i64 1970324836974593, i64 1970324836974611]
+// CHECK: @.offload_mapnames = private constant [9 x ptr] [ptr @{{.*}}, ptr @{{.*}}, ptr @{{.*}}, ptr @{{.*}}, ptr @{{.*}}, ptr @{{.*}}, ptr @{{.*}}, ptr @{{.*}}, ptr @{{.*}}]
+
+// CHECK: define void @_QQmain()
+
+// CHECK: %[[ARR_SECT_ALLOCA:.*]] = alloca { ptr, i64, i32, i8, i8, i8, i8, [1 x [3 x i64]] }, align 8
+// CHECK: %[[FULL_ARR_ALLOCA:.*]] = alloca { ptr, i64, i32, i8, i8, i8, i8, [1 x [3 x i64]] }, align 8
+// CHECK: %[[SCALAR_ALLOCA:.*]] = alloca { ptr, i64, i32, i8, i8, i8, i8 }, i64 1, align 8
+// CHECK: %[[FULL_ARR_SIZE6:.*]] = getelementptr { ptr, i64, i32, i8, i8, i8, i8, [1 x [3 x i64]] }, ptr %[[FULL_ARR_ALLOCA]], i32 0, i32 7, i64 0, i32 1
+// CHECK: %[[FULL_ARR_SIZE5:.*]] = load i64, ptr %[[FULL_ARR_SIZE6]], align 4
+// CHECK: %[[FULL_ARR_SIZE4:.*]] = sub i64 %[[FULL_ARR_SIZE5]], 1
+// CHECK: %[[ARR_SECT_OFFSET4:.*]] = getelementptr { ptr, i64, i32, i8, i8, i8, i8, [1 x [3 x i64]] }, ptr %[[ARR_SECT_ALLOCA]], i32 0, i32 7, i64 0, i32 0
+// CHECK: %[[ARR_SECT_OFFSET3:.*]] = load i64, ptr %[[ARR_SECT_OFFSET4]], align 4
+// CHECK: %[[ARR_SECT_OFFSET2:.*]] = sub i64 2, %[[ARR_SECT_OFFSET3]]
+// CHECK: %[[ARR_SECT_SIZE4:.*]] = sub i64 5, %[[ARR_SECT_OFFSET3]]
+// CHECK: %[[SCALAR_BASE:.*]] = getelementptr { ptr, i64, i32, i8, i8, i8, i8 }, ptr %[[SCALAR_ALLOCA]], i32 0, i32 0
+// CHECK: %[[FULL_ARR_SIZE3:.*]] = sub i64 %[[FULL_ARR_SIZE4]], 0
+// CHECK: %[[FULL_ARR_SIZE2:.*]] = add i64 %[[FULL_ARR_SIZE3]], 1
+// CHECK: %[[FULL_ARR_SIZE1:.*]] = mul i64 1, %[[FULL_ARR_SIZE2]]
+// CHECK: %[[FULL_ARR_SIZE:.*]] = mul i64 %[[FULL_ARR_SIZE1]], 4
+// CHECK: %[[ARR_SECT_SIZE3:.*]] = sub i64 %[[ARR_SECT_SIZE4]], %[[ARR_SECT_OFFSET2]]
+// CHECK: %[[ARR_SECT_SIZE2:.*]] = add i64 %[[ARR_SECT_SIZE3]], 1
+// CHECK: %[[ARR_SECT_SIZE1:.*]] = mul i64 1, %[[ARR_SECT_SIZE2]]
+// CHECK: %[[ARR_SECT_SIZE:.*]] = mul i64 %[[ARR_SECT_SIZE1]], 4
+// CHECK: %[[FULL_ARR_DESC_SIZE:.*]] = sdiv exact i64 sub (i64 ptrtoint (ptr getelementptr inbounds ({ ptr, i64, i32, i8, i8, i8, i8, [1 x [3 x i64]] }, ptr @_QFEfull_arr, i32 1) to i64), i64 ptrtoint (ptr @_QFEfull_arr to i64)), ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
+// CHECK: %[[LFULL_ARR:.*]] = load ptr, ptr @_QFEfull_arr, align 8
+// CHECK: %[[FULL_ARR_PTR:.*]] = getelementptr inbounds float, ptr %[[LFULL_ARR]], i64 0
+// CHECK: %[[ARR_SECT_DESC_SIZE:.*]] = sdiv exact i64 sub (i64 ptrtoint (ptr getelementptr inbounds ({ ptr, i64, i32, i8, i8, i8, i8, [1 x [3 x i64]] }, ptr @_QFEsect_arr, i32 1) to i64), i64 ptrtoint (ptr @_QFEsect_arr to i64)), ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
+// CHECK: %[[ARR_SECT_OFFSET1:.*]] = mul i64 %[[ARR_SECT_OFFSET2]], 1
+// CHECK: %[[LARR_SECT:.*]] = load ptr, ptr @_QFEsect_arr, align 8
+// CHECK: %[[ARR_SECT_PTR:.*]] = getelementptr inbounds i32, ptr %[[LARR_SECT]], i64 %[[ARR_SECT_OFFSET1]]
+// CHECK: %[[SCALAR_DESC_SZ4:.*]] = getelementptr { ptr, i64, i32, i8, i8, i8, i8 }, ptr %[[SCALAR_ALLOCA]], i32 1
+// CHECK: %[[SCALAR_DESC_SZ3:.*]] = ptrtoint ptr %[[SCALAR_DESC_SZ4]] to i64
+// CHECK: %[[SCALAR_DESC_SZ2:.*]] = ptrtoint ptr %[[SCALAR_ALLOCA]] to i64
+// CHECK: %[[SCALAR_DESC_SZ1:.*]] = sub i64 %[[SCALAR_DESC_SZ3]], %[[SCALAR_DESC_SZ2]]
+// CHECK: %[[SCALAR_DESC_SZ:.*]] = sdiv exact i64 %[[SCALAR_DESC_SZ1]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
+// CHECK: %[[SCALAR_PTR_LOAD:.*]] = load ptr, ptr %[[SCALAR_BASE]], align 8
+// CHECK: %[[SCALAR_PTR:.*]] = getelementptr inbounds float, ptr %[[SCALAR_PTR_LOAD]], i64 0
+
+// CHECK: %[[OFFLOADBASEPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
+// CHECK: store ptr @_QFEfull_arr, ptr %[[OFFLOADBASEPTRS]], align 8
+// CHECK: %[[OFFLOADPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_ptrs, i32 0, i32 0
+// CHECK: store ptr @_QFEfull_arr, ptr %[[OFFLOADPTRS]], align 8
+
+// CHECK: %[[OFFLOADSIZES:.*]] = getelementptr inbounds [9 x i64], ptr %.offload_sizes, i32 0, i32 0
+// CHECK: store i64 %[[FULL_ARR_DESC_SIZE]], ptr %[[OFFLOADSIZES]], align 8
+
+// CHECK: %[[OFFLOADBASEPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_baseptrs, i32 0, i32 1
+// CHECK: store ptr @_QFEfull_arr, ptr %[[OFFLOADBASEPTRS]], align 8
+// CHECK: %[[OFFLOADPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_ptrs, i32 0, i32 1
+// CHECK: store ptr @_QFEfull_arr, ptr %[[OFFLOADPTRS]], align 8
+// CHECK: %[[OFFLOADBASEPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_baseptrs, i32 0, i32 2
+// CHECK: store ptr @_QFEfull_arr, ptr %[[OFFLOADBASEPTRS]], align 8
+// CHECK: %[[OFFLOADPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_ptrs, i32 0, i32 2
+// CHECK: store ptr %[[FULL_ARR_PTR]], ptr %[[OFFLOADPTRS]], align 8
+// CHECK: %[[OFFLOADSIZES:.*]] = getelementptr inbounds [9 x i64], ptr %.offload_sizes, i32 0, i32 2
+// CHECK: store i64 %[[FULL_ARR_SIZE]], ptr %[[OFFLOADSIZES]], align 8
+
+// CHECK: %[[OFFLOADBASEPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_baseptrs, i32 0, i32 3
+// CHECK: store ptr @_QFEsect_arr, ptr %[[OFFLOADBASEPTRS]], align 8
+// CHECK: %[[OFFLOADPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_ptrs, i32 0, i32 3
+// CHECK: store ptr @_QFEsect_arr, ptr %[[OFFLOADPTRS]], align 8
+// CHECK: %[[OFFLOADSIZES:.*]] = getelementptr inbounds [9 x i64], ptr %.offload_sizes, i32 0, i32 3
+// CHECK: store i64 %[[ARR_SECT_DESC_SIZE]], ptr %[[OFFLOADSIZES]], align 8
+// CHECK: %[[OFFLOADBASEPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_baseptrs, i32 0, i32 4
+// CHECK: store ptr @_QFEsect_arr, ptr %[[OFFLOADBASEPTRS]], align 8
+// CHECK: %[[OFFLOADPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_ptrs, i32 0, i32 4
+// CHECK: store ptr @_QFEsect_arr, ptr %[[OFFLOADPTRS]], align 8
+// CHECK: %[[OFFLOADBASEPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_baseptrs, i32 0, i32 5
+// CHECK: store ptr @_QFEsect_arr, ptr %[[OFFLOADBASEPTRS]], align 8
+// CHECK: %[[OFFLOADPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_ptrs, i32 0, i32 5
+// CHECK: store ptr %[[ARR_SECT_PTR]], ptr %[[OFFLOADPTRS]], align 8
+// CHECK: %[[OFFLOADSIZES:.*]] = getelementptr inbounds [9 x i64], ptr %.offload_sizes, i32 0, i32 5
+// CHECK: store i64 %[[ARR_SECT_SIZE]], ptr %[[OFFLOADSIZES]], align 8
+
+// CHECK: %[[OFFLOADBASEPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_baseptrs, i32 0, i32 6
+// CHECK: store ptr %[[SCALAR_ALLOCA]], ptr %[[OFFLOADBASEPTRS]], align 8
+// CHECK: %[[OFFLOADPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_ptrs, i32 0, i32 6
+// CHECK: store ptr %[[SCALAR_ALLOCA]], ptr %[[OFFLOADPTRS]], align 8
+// CHECK: %[[OFFLOADSIZES:.*]] = getelementptr inbounds [9 x i64], ptr %.offload_sizes, i32 0, i32 6
+// CHECK: store i64 %[[SCALAR_DESC_SZ]], ptr %[[OFFLOADSIZES]], align 8
+// CHECK: %[[OFFLOADBASEPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_baseptrs, i32 0, i32 7
+// CHECK: store ptr %[[SCALAR_ALLOCA]], ptr %[[OFFLOADBASEPTRS]], align 8
+// CHECK: %[[OFFLOADPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_ptrs, i32 0, i32 7
+// CHECK: store ptr %[[SCALAR_ALLOCA]], ptr %[[OFFLOADPTRS]], align 8
+// CHECK: %[[OFFLOADBASEPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_baseptrs, i32 0, i32 8
+// CHECK: store ptr %[[SCALAR_BASE]], ptr %[[OFFLOADBASEPTRS]], align 8
+// CHECK: %[[OFFLOADPTRS:.*]] = getelementptr inbounds [9 x ptr], ptr %.offload_ptrs, i32 0, i32 8
+// CHECK: store ptr %[[SCALAR_PTR]], ptr %[[OFFLOADPTRS]], align 8
diff --git a/openmp/libomptarget/test/offloading/fortran/allocatable-array-section-1d-bounds.f90 b/openmp/libomptarget/test/offloading/fortran/allocatable-array-section-1d-bounds.f90
new file mode 100644
index 00000000000000..99dbe99d40497c
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/fortran/allocatable-array-section-1d-bounds.f90
@@ -0,0 +1,46 @@
+! Offloading test checking interaction of a
+! two 1-D allocatable arrays with a target region
+! while providing the map upper and lower bounds
+! REQUIRES: flang, amdgcn-amd-amdhsa
+! UNSUPPORTED: nvptx64-nvidia-cuda
+! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+! UNSUPPORTED: aarch64-unknown-linux-gnu
+! UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+! UNSUPPORTED: x86_64-pc-linux-gnu
+! UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+program main
+ integer, allocatable :: sp_read(:), sp_write(:)
+ allocate(sp_read(10))
+ allocate(sp_write(10))
+
+ do i = 1, 10
+ sp_read(i) = i
+ sp_write(i) = 0
+ end do
+
+ !$omp target map(tofrom:sp_read(2:6)) map(tofrom:sp_write(2:6))
+ do i = 1, 10
+ sp_write(i) = sp_read(i)
+ end do
+ !$omp end target
+
+ do i = 1, 10
+ print *, sp_write(i)
+ end do
+
+ deallocate(sp_read)
+ deallocate(sp_write)
+end program
+
+! CHECK: 0
+! CHECK: 2
+! CHECK: 3
+! CHECK: 4
+! CHECK: 5
+! CHECK: 6
+! CHECK: 0
+! CHECK: 0
+! CHECK: 0
+! CHECK: 0
diff --git a/openmp/libomptarget/test/offloading/fortran/allocatable-array-section-3d-bounds.f90 b/openmp/libomptarget/test/offloading/fortran/allocatable-array-section-3d-bounds.f90
new file mode 100644
index 00000000000000..0786e0fd744e78
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/fortran/allocatable-array-section-3d-bounds.f90
@@ -0,0 +1,44 @@
+! Offloading test checking interaction of allocatables
+! with multi-dimensional bounds (3-D in this case) and
+! a target region
+! REQUIRES: flang, amdgcn-amd-amdhsa
+! UNSUPPORTED: nvptx64-nvidia-cuda
+! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+! UNSUPPORTED: aarch64-unknown-linux-gnu
+! UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+! UNSUPPORTED: x86_64-pc-linux-gnu
+! UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+program main
+ integer, allocatable :: inArray(:,:,:)
+ integer, allocatable :: outArray(:,:,:)
+
+ allocate(inArray(3,3,3))
+ allocate(outArray(3,3,3))
+
+ do i = 1, 3
+ do j = 1, 3
+ do k = 1, 3
+ inArray(i, j, k) = 42
+ outArray(i, j, k) = 0
+ end do
+ end do
+ end do
+
+!$omp target map(tofrom:inArray(1:3, 1:3, 2:2), outArray(1:3, 1:3, 1:3))
+ do j = 1, 3
+ do k = 1, 3
+ outArray(k, j, 2) = inArray(k, j, 2)
+ end do
+ end do
+!$omp end target
+
+print *, outArray
+
+deallocate(inArray)
+deallocate(outArray)
+
+end program
+
+! CHECK: 0 0 0 0 0 0 0 0 0 42 42 42 42 42 42 42 42 42 0 0 0 0 0 0 0 0 0
diff --git a/openmp/libomptarget/test/offloading/fortran/allocatable-map-scopes.f95 b/openmp/libomptarget/test/offloading/fortran/allocatable-map-scopes.f95
new file mode 100644
index 00000000000000..bb47d3de96d2a9
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/fortran/allocatable-map-scopes.f95
@@ -0,0 +1,66 @@
+! Offloading test checking interaction of allocatables
+! with target in different scopes
+! REQUIRES: flang, amdgcn-amd-amdhsa
+! UNSUPPORTED: nvptx64-nvidia-cuda
+! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+! UNSUPPORTED: aarch64-unknown-linux-gnu
+! UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+! UNSUPPORTED: x86_64-pc-linux-gnu
+! UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+module test
+ contains
+ subroutine func_arg(arg_alloc)
+ integer, allocatable, intent (inout) :: arg_alloc(:)
+
+ !$omp target map(tofrom: arg_alloc)
+ do index = 1, 10
+ arg_alloc(index) = arg_alloc(index) + index
+ end do
+ !$omp end target
+
+ print *, arg_alloc
+ end subroutine func_arg
+end module
+
+subroutine func
+ integer, allocatable :: local_alloc(:)
+ allocate(local_alloc(10))
+
+ !$omp target map(tofrom: local_alloc)
+ do index = 1, 10
+ local_alloc(index) = index
+ end do
+ !$omp end target
+
+ print *, local_alloc
+
+ deallocate(local_alloc)
+end subroutine func
+
+
+program main
+ use test
+ integer, allocatable :: map_ptr(:)
+
+ allocate(map_ptr(10))
+
+ !$omp target map(tofrom: map_ptr)
+ do index = 1, 10
+ map_ptr(index) = index
+ end do
+ !$omp end target
+
+ call func
+
+ print *, map_ptr
+
+ call func_arg(map_ptr)
+
+ deallocate(map_ptr)
+end program
+
+! CHECK: 1 2 3 4 5 6 7 8 9 10
+! CHECK: 1 2 3 4 5 6 7 8 9 10
+! CHECK: 2 4 6 8 10 12 14 16 18 20
diff --git a/openmp/libomptarget/test/offloading/fortran/failing/target_single_value_allocate.f90 b/openmp/libomptarget/test/offloading/fortran/failing/target_single_value_allocate.f90
new file mode 100644
index 00000000000000..adb5bad4278db6
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/fortran/failing/target_single_value_allocate.f90
@@ -0,0 +1,27 @@
+! Offloading test checking interaction of a
+! non-array allocatable with a target region
+! REQUIRES: flang, amdgcn-amd-amdhsa
+! UNSUPPORTED: nvptx64-nvidia-cuda
+! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+! UNSUPPORTED: aarch64-unknown-linux-gnu
+! UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+! UNSUPPORTED: x86_64-pc-linux-gnu
+! UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+! XFAIL: *
+program main
+ integer, allocatable :: test
+ allocate(test)
+ test = 10
+
+!$omp target map(tofrom:test)
+ test = 50
+!$omp end target
+
+ print *, test
+
+ deallocate(test)
+end program
+
+! CHECK:
diff --git a/openmp/libomptarget/test/offloading/fortran/pointer-scopes-enter-exit-map.f90 b/openmp/libomptarget/test/offloading/fortran/pointer-scopes-enter-exit-map.f90
new file mode 100644
index 00000000000000..dee75af06927be
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/fortran/pointer-scopes-enter-exit-map.f90
@@ -0,0 +1,83 @@
+! Offloading test checking interaction of pointers
+! with target in different scopes
+! REQUIRES: flang, amdgcn-amd-amdhsa
+! UNSUPPORTED: nvptx64-nvidia-cuda
+! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+! UNSUPPORTED: aarch64-unknown-linux-gnu
+! UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+! UNSUPPORTED: x86_64-pc-linux-gnu
+! UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+module test
+ contains
+ subroutine func_arg(arg_alloc)
+ integer, pointer, intent (inout) :: arg_alloc(:)
+
+ !$omp target enter data map(alloc: arg_alloc)
+
+ !$omp target
+ do index = 1, 10
+ arg_alloc(index) = arg_alloc(index) + index
+ end do
+ !$omp end target
+
+ !$omp target exit data map(from: arg_alloc)
+
+ !$omp target exit data map(delete: arg_alloc)
+
+ print *, arg_alloc
+ end subroutine func_arg
+end module
+
+subroutine func
+ integer, pointer :: local_alloc(:)
+ allocate(local_alloc(10))
+
+ !$omp target enter data map(alloc: local_alloc)
+
+ !$omp target
+ do index = 1, 10
+ local_alloc(index) = index
+ end do
+ !$omp end target
+
+ !$omp target exit data map(from: local_alloc)
+
+ !$omp target exit data map(delete: local_alloc)
+
+ print *, local_alloc
+
+ deallocate(local_alloc)
+end subroutine func
+
+
+program main
+ use test
+ integer, pointer :: map_ptr(:)
+ allocate(map_ptr(10))
+
+ !$omp target enter data map(alloc: map_ptr)
+
+ !$omp target
+ do index = 1, 10
+ map_ptr(index) = index
+ end do
+ !$omp end target
+
+ !$omp target exit data map(from: map_ptr)
+
+ !$omp target exit data map(delete: map_ptr)
+
+ call func
+
+ print *, map_ptr
+
+ call func_arg(map_ptr)
+
+ deallocate(map_ptr)
+end program
+
+! CHECK: 1 2 3 4 5 6 7 8 9 10
+! CHECK: 1 2 3 4 5 6 7 8 9 10
+! CHECK: 2 4 6 8 10 12 14 16 18 20
diff --git a/openmp/libomptarget/test/offloading/fortran/pointer-target-array-section-3d-bounds.f90 b/openmp/libomptarget/test/offloading/fortran/pointer-target-array-section-3d-bounds.f90
new file mode 100644
index 00000000000000..ff2298cf5dbc9d
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/fortran/pointer-target-array-section-3d-bounds.f90
@@ -0,0 +1,43 @@
+! Offloading test checking interaction of pointer
+! and target with target where 3-D bounds have
+! been specified
+! REQUIRES: flang, amdgcn-amd-amdhsa
+! UNSUPPORTED: nvptx64-nvidia-cuda
+! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+! UNSUPPORTED: aarch64-unknown-linux-gnu
+! UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+! UNSUPPORTED: x86_64-pc-linux-gnu
+! UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+program main
+ integer, pointer :: inArray(:,:,:)
+ integer, pointer :: outArray(:,:,:)
+ integer, target :: in(3,3,3)
+ integer, target :: out(3,3,3)
+
+ inArray => in
+ outArray => out
+
+ do i = 1, 3
+ do j = 1, 3
+ do k = 1, 3
+ inArray(i, j, k) = 42
+ outArray(i, j, k) = 0
+ end do
+ end do
+ end do
+
+!$omp target map(tofrom:inArray(1:3, 1:3, 2:2), outArray(1:3, 1:3, 1:3))
+ do j = 1, 3
+ do k = 1, 3
+ outArray(k, j, 2) = inArray(k, j, 2)
+ end do
+ end do
+!$omp end target
+
+ print *, outArray
+
+end program
+
+! CHECK: 0 0 0 0 0 0 0 0 0 42 42 42 42 42 42 42 42 42 0 0 0 0 0 0 0 0 0
diff --git a/openmp/libomptarget/test/offloading/fortran/pointer-target-map-scopes.f95 b/openmp/libomptarget/test/offloading/fortran/pointer-target-map-scopes.f95
new file mode 100644
index 00000000000000..d9a7000719f0e8
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/fortran/pointer-target-map-scopes.f95
@@ -0,0 +1,64 @@
+! Offloading test checking interaction of pointer
+! and target with target across multiple scopes
+! REQUIRES: flang, amdgcn-amd-amdhsa
+! UNSUPPORTED: nvptx64-nvidia-cuda
+! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+! UNSUPPORTED: aarch64-unknown-linux-gnu
+! UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+! UNSUPPORTED: x86_64-pc-linux-gnu
+! UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+module test
+ contains
+ subroutine func_arg(arg_alloc)
+ integer, pointer, intent (inout) :: arg_alloc(:)
+
+ !$omp target map(tofrom: arg_alloc)
+ do index = 1, 10
+ arg_alloc(index) = arg_alloc(index) + index
+ end do
+ !$omp end target
+
+ print *, arg_alloc
+ end subroutine func_arg
+end module
+
+subroutine func
+ integer, pointer :: local_alloc(:)
+ integer, target :: b(10)
+ local_alloc => b
+
+ !$omp target map(tofrom: local_alloc)
+ do index = 1, 10
+ local_alloc(index) = index
+ end do
+ !$omp end target
+
+ print *, local_alloc
+ end subroutine func
+
+
+ program main
+ use test
+ integer, pointer :: map_ptr(:)
+ integer, target :: b(10)
+
+ map_ptr => b
+
+ !$omp target map(tofrom: map_ptr)
+ do index = 1, 10
+ map_ptr(index) = index
+ end do
+ !$omp end target
+
+ call func
+
+ print *, map_ptr
+
+ call func_arg(map_ptr)
+end program
+
+!CHECK: 1 2 3 4 5 6 7 8 9 10
+!CHECK: 1 2 3 4 5 6 7 8 9 10
+!CHECK: 2 4 6 8 10 12 14 16 18 20
diff --git a/openmp/libomptarget/test/offloading/fortran/target_enter_exit_allocatables.f90 b/openmp/libomptarget/test/offloading/fortran/target_enter_exit_allocatables.f90
new file mode 100644
index 00000000000000..865be95ba9682d
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/fortran/target_enter_exit_allocatables.f90
@@ -0,0 +1,44 @@
+! Offloading test checking interaction of allocatables
+! with enter, exit and target
+! REQUIRES: flang, amdgcn-amd-amdhsa
+! UNSUPPORTED: nvptx64-nvidia-cuda
+! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+! UNSUPPORTED: aarch64-unknown-linux-gnu
+! UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+! UNSUPPORTED: x86_64-pc-linux-gnu
+! UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+program main
+ integer, allocatable :: A(:)
+ allocate(A(10))
+
+ !$omp target enter data map(alloc: A)
+
+ !$omp target
+ do I = 1, 10
+ A(I) = I
+ end do
+ !$omp end target
+
+ !$omp target exit data map(from: A)
+
+ !$omp target exit data map(delete: A)
+
+ do i = 1, 10
+ print *, A(i)
+ end do
+
+ deallocate(A)
+end program
+
+! CHECK: 1
+! CHECK: 2
+! CHECK: 3
+! CHECK: 4
+! CHECK: 5
+! CHECK: 6
+! CHECK: 7
+! CHECK: 8
+! CHECK: 9
+! CHECK: 10
diff --git a/openmp/libomptarget/test/offloading/fortran/target_enter_exit_array.f90 b/openmp/libomptarget/test/offloading/fortran/target_enter_exit_array.f90
new file mode 100644
index 00000000000000..4a9fb6ee177f65
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/fortran/target_enter_exit_array.f90
@@ -0,0 +1,41 @@
+! Offloading test checking interaction of fixed size
+! arrays with enter, exit and target
+! REQUIRES: flang, amdgcn-amd-amdhsa
+! UNSUPPORTED: nvptx64-nvidia-cuda
+! UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+! UNSUPPORTED: aarch64-unknown-linux-gnu
+! UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+! UNSUPPORTED: x86_64-pc-linux-gnu
+! UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+! RUN: %libomptarget-compile-fortran-run-and-check-generic
+program main
+ integer :: A(10)
+
+ !$omp target enter data map(alloc: A)
+
+ !$omp target
+ do I = 1, 10
+ A(I) = I
+ end do
+ !$omp end target
+
+ !$omp target exit data map(from: A)
+
+ !$omp target exit data map(delete: A)
+
+ do i = 1, 10
+ print *, A(i)
+ end do
+end program
+
+! CHECK: 1
+! CHECK: 2
+! CHECK: 3
+! CHECK: 4
+! CHECK: 5
+! CHECK: 6
+! CHECK: 7
+! CHECK: 8
+! CHECK: 9
+! CHECK: 10
>From c53a91b5069dda02d71ede0a55f259db75e42420 Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Mon, 4 Dec 2023 08:18:51 -0600
Subject: [PATCH 02/19] missed rebase artifact
---
flang/test/Lower/OpenMP/FIR/array-bounds.f90 | 1 -
1 file changed, 1 deletion(-)
diff --git a/flang/test/Lower/OpenMP/FIR/array-bounds.f90 b/flang/test/Lower/OpenMP/FIR/array-bounds.f90
index 7b9802585f9425..ec0ecd3c90bb9e 100644
--- a/flang/test/Lower/OpenMP/FIR/array-bounds.f90
+++ b/flang/test/Lower/OpenMP/FIR/array-bounds.f90
@@ -63,7 +63,6 @@ end subroutine assumed_shape_array
!ALL-LABEL: func.func @_QMassumed_array_routinesPassumed_size_array(
!ALL-SAME: %[[ARG0:.*]]: !fir.ref<!fir.array<?xi32>> {fir.bindc_name = "arr_read_write"})
-!ALL: %[[UNDEF:.*]] = fir.undefined index
!ALL: %[[ALLOCA:.*]] = fir.alloca i32 {bindc_name = "i", uniq_name = "_QMassumed_array_routinesFassumed_size_arrayEi"}
!ALL: %[[C0:.*]] = arith.constant 1 : index
!ALL: %[[C1:.*]] = arith.constant 1 : index
>From 2f3176cdadb553afa585b2577e03a54038e4ea60 Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Mon, 4 Dec 2023 15:08:01 -0600
Subject: [PATCH 03/19] Remove unneccesary blocking for implicit cases
This was intended to allow for enter + exit data to work in conjunction
with descriptor based type mapping, but after further testing it seems
unneccesary and actually hinders the regular implicit capture case. It's
also best to keep this lowering as simple as possible for the case where
we're actually mapping a derived type/structure.
---
.../Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp | 14 +-------------
1 file changed, 1 insertion(+), 13 deletions(-)
diff --git a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
index 15152cf6a8357a..aea1a1bfedbcba 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
@@ -1989,19 +1989,7 @@ static void genMapInfos(llvm::IRBuilderBase &builder,
for (unsigned long i : primaryMapIdx) {
auto mapInfoOp = mlir::dyn_cast<mlir::omp::MapInfoOp>(mapData.MapClause[i]);
- bool isImplicit =
- mapInfoOp.getMapType().value() &
- static_cast<
- std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
- llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT);
-
- // TODO: Look into more thoroughly handling implicit semantics, this
- // mimics what Clang currently does with structures, which is map the
- // structure and not the internals i.e. the descriptor without the pointer
- // itself, so not a deep copy, which may be incorrect for
- // allocatables/pointers, but allows explicit mapping and initial
- // enter/exit handling
- if (!isImplicit && !mapInfoOp.getMembers().empty()) {
+ if (!mapInfoOp.getMembers().empty()) {
processMapWithMembersOf(moduleTranslation, builder, *ompBuilder, dl,
combinedInfo, mapData, i, isTargetParams);
continue;
>From b8fa58edfd4a7878cf20be13061acb5d1cf9ab99 Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Tue, 5 Dec 2023 09:48:32 -0600
Subject: [PATCH 04/19] Fix the implicit map types for certain box captures
---
flang/lib/Lower/OpenMP.cpp | 18 ++++++++++--------
1 file changed, 10 insertions(+), 8 deletions(-)
diff --git a/flang/lib/Lower/OpenMP.cpp b/flang/lib/Lower/OpenMP.cpp
index ae4e295f59625f..d0838caaa1243d 100644
--- a/flang/lib/Lower/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP.cpp
@@ -2813,14 +2813,16 @@ genTargetOp(Fortran::lower::AbstractConverter &converter,
llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT;
mlir::omp::VariableCaptureKind captureKind =
mlir::omp::VariableCaptureKind::ByRef;
- if (auto refType = baseOp.getType().dyn_cast<fir::ReferenceType>()) {
- auto eleType = refType.getElementType();
- if (fir::isa_trivial(eleType) || fir::isa_char(eleType)) {
- captureKind = mlir::omp::VariableCaptureKind::ByCopy;
- } else if (!fir::isa_builtin_cptr_type(eleType)) {
- mapFlag |= llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_TO;
- mapFlag |= llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_FROM;
- }
+
+ mlir::Type eleType = baseOp.getType();
+ if (auto refType = baseOp.getType().dyn_cast<fir::ReferenceType>())
+ eleType = refType.getElementType();
+
+ if (fir::isa_trivial(eleType) || fir::isa_char(eleType)) {
+ captureKind = mlir::omp::VariableCaptureKind::ByCopy;
+ } else if (!fir::isa_builtin_cptr_type(eleType)) {
+ mapFlag |= llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_TO;
+ mapFlag |= llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_FROM;
}
mlir::Value mapOp;
>From 6e12ec1b805510f1a4f857d633204d0599b90c2e Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Tue, 5 Dec 2023 11:05:34 -0600
Subject: [PATCH 05/19] Forgotten commit of minor test adjustments
---
flang/test/Integration/OpenMP/map-types-and-sizes.f90 | 8 ++++----
1 file changed, 4 insertions(+), 4 deletions(-)
diff --git a/flang/test/Integration/OpenMP/map-types-and-sizes.f90 b/flang/test/Integration/OpenMP/map-types-and-sizes.f90
index 97723aae57d57c..35f1dd8bcaea60 100644
--- a/flang/test/Integration/OpenMP/map-types-and-sizes.f90
+++ b/flang/test/Integration/OpenMP/map-types-and-sizes.f90
@@ -30,8 +30,8 @@ subroutine mapType_array
!$omp end target
end subroutine mapType_array
-!CHECK: @.offload_sizes{{.*}} = private unnamed_addr constant [1 x i64] [i64 24]
-!CHECK: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 547]
+!CHECK: @.offload_sizes{{.*}} = private unnamed_addr constant [3 x i64] [i64 0, i64 24, i64 4]
+!CHECK: @.offload_maptypes{{.*}} = private unnamed_addr constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976711187]
subroutine mapType_ptr
integer, pointer :: a
!$omp target
@@ -39,8 +39,8 @@ subroutine mapType_ptr
!$omp end target
end subroutine mapType_ptr
-!CHECK: @.offload_sizes{{.*}} = private unnamed_addr constant [1 x i64] [i64 24]
-!CHECK: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 547]
+!CHECK: @.offload_sizes{{.*}} = private unnamed_addr constant [3 x i64] [i64 0, i64 24, i64 4]
+!CHECK: @.offload_maptypes{{.*}} = private unnamed_addr constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976711187]
subroutine mapType_allocatable
integer, allocatable :: a
allocate(a)
>From 43bf4da35d05cac19c3eb87a5c66a792d83f310b Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Tue, 5 Dec 2023 13:45:26 -0600
Subject: [PATCH 06/19] Add fir::isTypeWithDescriptor function to FIRType
---
flang/include/flang/Optimizer/Dialect/FIRType.h | 3 +++
flang/lib/Lower/OpenMP.cpp | 8 ++------
flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp | 4 +---
flang/lib/Optimizer/Dialect/FIRType.cpp | 7 +++++++
4 files changed, 13 insertions(+), 9 deletions(-)
diff --git a/flang/include/flang/Optimizer/Dialect/FIRType.h b/flang/include/flang/Optimizer/Dialect/FIRType.h
index ecfa9839617dab..01ce0517d89b60 100644
--- a/flang/include/flang/Optimizer/Dialect/FIRType.h
+++ b/flang/include/flang/Optimizer/Dialect/FIRType.h
@@ -313,6 +313,9 @@ bool isBoxNone(mlir::Type ty);
/// e.g. !fir.box<!fir.type<derived>>
bool isBoxedRecordType(mlir::Type ty);
+/// Return true iff `ty` is a type that contains descriptor information.
+bool isTypeWithDescriptor(mlir::Type ty);
+
/// Return true iff `ty` is a scalar boxed record type.
/// e.g. !fir.box<!fir.type<derived>>
/// !fir.box<!fir.heap<!fir.type<derived>>>
diff --git a/flang/lib/Lower/OpenMP.cpp b/flang/lib/Lower/OpenMP.cpp
index d0838caaa1243d..e07ffa90406a85 100644
--- a/flang/lib/Lower/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP.cpp
@@ -1862,9 +1862,7 @@ bool ClauseProcessor::processMap(
auto origSymbol =
converter.getSymbolAddress(*getOmpObjectSymbol(ompObject));
mlir::Value mapOp, symAddr;
- if (fir::isPointerType(origSymbol.getType()) ||
- fir::isAllocatableType(origSymbol.getType()) ||
- fir::isAssumedShape(origSymbol.getType())) {
+ if (fir::isTypeWithDescriptor(origSymbol.getType())) {
symAddr = origSymbol;
mapOp = processDescriptorTypeMappings(
semanticsContext, stmtCtx, converter, clauseLocation,
@@ -2826,9 +2824,7 @@ genTargetOp(Fortran::lower::AbstractConverter &converter,
}
mlir::Value mapOp;
- if (fir::isPointerType(baseOp.getType()) ||
- fir::isAllocatableType(baseOp.getType()) ||
- fir::isAssumedShape(baseOp.getType())) {
+ if (fir::isTypeWithDescriptor(baseOp.getType())) {
mapOp = processDescriptorTypeMappings(
semanticsContext, stmtCtx, converter, baseOp.getLoc(), baseOp,
baseAddr, bounds, name.str(), mapFlag);
diff --git a/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp b/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
index 5de95ed807b3fb..3dadfe74817a45 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
@@ -59,9 +59,7 @@ struct MapInfoOpConversion
for (mlir::NamedAttribute attr : curOp->getAttrs()) {
if (auto typeAttr = mlir::dyn_cast<mlir::TypeAttr>(attr.getValue())) {
mlir::Type newAttr;
- if (fir::isPointerType(typeAttr.getValue()) ||
- fir::isAllocatableType(typeAttr.getValue()) ||
- fir::isAssumedShape(typeAttr.getValue())) {
+ if (fir::isTypeWithDescriptor(typeAttr.getValue())) {
newAttr = lowerTy().convertBoxTypeAsStruct(
mlir::cast<fir::BaseBoxType>(typeAttr.getValue()));
} else {
diff --git a/flang/lib/Optimizer/Dialect/FIRType.cpp b/flang/lib/Optimizer/Dialect/FIRType.cpp
index d0c7bae674b6cf..9323a019dd7b55 100644
--- a/flang/lib/Optimizer/Dialect/FIRType.cpp
+++ b/flang/lib/Optimizer/Dialect/FIRType.cpp
@@ -330,6 +330,13 @@ bool isAllocatableOrPointerArray(mlir::Type ty) {
return false;
}
+bool isTypeWithDescriptor(mlir::Type ty) {
+ if (fir::isPointerType(ty) || fir::isAllocatableType(ty) ||
+ fir::isAssumedShape(ty))
+ return true;
+ return false;
+}
+
bool isPolymorphicType(mlir::Type ty) {
if (auto refTy = fir::dyn_cast_ptrEleTy(ty))
ty = refTy;
>From 8480af33849707273a091aae317bdc23466ccfd0 Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Tue, 5 Dec 2023 13:47:25 -0600
Subject: [PATCH 07/19] remove addendum mapping for now and add TODO
---
flang/lib/Lower/OpenMP.cpp | 18 ++----------------
1 file changed, 2 insertions(+), 16 deletions(-)
diff --git a/flang/lib/Lower/OpenMP.cpp b/flang/lib/Lower/OpenMP.cpp
index e07ffa90406a85..8166ab00480b41 100644
--- a/flang/lib/Lower/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP.cpp
@@ -1769,22 +1769,8 @@ static mlir::omp::MapInfoOp processDescriptorTypeMappings(
mapCaptureType),
mlir::omp::VariableCaptureKind::ByRef, descDataBaseAddr.getType()));
- if (fir::BaseBoxType baseBoxTy = llvm::dyn_cast<fir::BaseBoxType>(
- fir::unwrapRefType(descriptor.getType()))) {
- if (fir::boxHasAddendum(baseBoxTy)) {
- mlir::Value addendumAddrAddr = firOpBuilder.create<fir::BoxOffsetOp>(
- loc, descriptor, fir::BoxFieldAttr::derived_type);
- // NOTE: We have no addendumAddr, but we must provide a varPtr alongside
- // the varPtrPtr so we replicate it across the two for now.
- descriptorBaseAddrMembers.push_back(createMapInfoOp(
- firOpBuilder, loc, addendumAddrAddr, addendumAddrAddr, asFortran, {},
- {},
- static_cast<
- std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
- mapCaptureType),
- mlir::omp::VariableCaptureKind::ByRef, addendumAddrAddr.getType()));
- }
- }
+ // TODO: map the addendum segment of the descriptor, similarly to the abose
+ // base address/data pointer member.
return createMapInfoOp(
firOpBuilder, loc, descriptor, mlir::Value{}, asFortran, {},
>From 77d27b4d2a390f5da41da9786a755f6fc63dc8c5 Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Tue, 5 Dec 2023 14:06:59 -0600
Subject: [PATCH 08/19] Add comment for AssumedShape case for further context
---
flang/lib/Lower/OpenMP.cpp | 8 ++++++++
1 file changed, 8 insertions(+)
diff --git a/flang/lib/Lower/OpenMP.cpp b/flang/lib/Lower/OpenMP.cpp
index 8166ab00480b41..7743b2ec150f23 100644
--- a/flang/lib/Lower/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP.cpp
@@ -1751,6 +1751,14 @@ static mlir::omp::MapInfoOp processDescriptorTypeMappings(
fir::FirOpBuilder &firOpBuilder = converter.getFirOpBuilder();
mlir::Value descriptor = descriptorAddr;
+
+ // The fir::BoxOffsetOp only works with !fir.ref<!fir.box<...>> types, as
+ // allowing it to access non-reference box operations can cause some
+ // problematic SSA IR. However, in the case of assumed shape's the type
+ // is not a !fir.ref, in these cases to retrieve the appropriate
+ // !fir.ref<!fir.box<...>> to access the data we need to map we must
+ // perform an alloca and then store to it and retrieve the data from the new
+ // alloca.
if (fir::isAssumedShape(fir::unwrapRefType(descriptorAddr.getType()))) {
mlir::OpBuilder::InsertPoint insPt = firOpBuilder.saveInsertionPoint();
firOpBuilder.setInsertionPointToStart(firOpBuilder.getAllocaBlock());
>From 7bcd5380877e5faf6e9498da302d3921fdd7340f Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Tue, 5 Dec 2023 14:30:10 -0600
Subject: [PATCH 09/19] Change llvm::SmallVector<mlir::Value> ->
mlir::ValueRange
---
flang/lib/Lower/OpenMP.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/flang/lib/Lower/OpenMP.cpp b/flang/lib/Lower/OpenMP.cpp
index 7743b2ec150f23..f6c78aec30edcd 100644
--- a/flang/lib/Lower/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP.cpp
@@ -1745,7 +1745,7 @@ static mlir::omp::MapInfoOp processDescriptorTypeMappings(
Fortran::lower::StatementContext &stmtCtx,
Fortran::lower::AbstractConverter &converter, mlir::Location loc,
mlir::Value descriptorAddr, mlir::Value descDataBaseAddr,
- mlir::SmallVector<mlir::Value> &bounds, std::string asFortran,
+ mlir::ValueRange bounds, std::string asFortran,
llvm::omp::OpenMPOffloadMappingFlags mapCaptureType) {
llvm::SmallVector<mlir::Value> descriptorBaseAddrMembers;
fir::FirOpBuilder &firOpBuilder = converter.getFirOpBuilder();
>From 9cdf3188ebdaa244e0ebf02d8ef899953c2cc618 Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Tue, 5 Dec 2023 14:47:59 -0600
Subject: [PATCH 10/19] Refactor things into an CodeGenOpenMP header
---
.../include/flang/Optimizer/CodeGen/CodeGen.h | 7 -----
.../flang/Optimizer/CodeGen/CodeGenOpenMP.h | 27 +++++++++++++++++++
flang/lib/Optimizer/CodeGen/CodeGen.cpp | 1 +
flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp | 3 ++-
4 files changed, 30 insertions(+), 8 deletions(-)
create mode 100644 flang/include/flang/Optimizer/CodeGen/CodeGenOpenMP.h
diff --git a/flang/include/flang/Optimizer/CodeGen/CodeGen.h b/flang/include/flang/Optimizer/CodeGen/CodeGen.h
index 27bea240fa8e8e..5ea96c900bc630 100644
--- a/flang/include/flang/Optimizer/CodeGen/CodeGen.h
+++ b/flang/include/flang/Optimizer/CodeGen/CodeGen.h
@@ -19,7 +19,6 @@
namespace fir {
struct NameUniquer;
-class LLVMTypeConverter;
#define GEN_PASS_DECL_FIRTOLLVMLOWERING
#define GEN_PASS_DECL_CODEGENREWRITE
@@ -81,12 +80,6 @@ std::unique_ptr<mlir::Pass> createLLVMDialectToLLVMPass(
std::unique_ptr<mlir::Pass> createBoxedProcedurePass();
std::unique_ptr<mlir::Pass> createBoxedProcedurePass(bool useThunks);
-/// Specialised conversion patterns of OpenMP operations for FIR to LLVM
-/// dialect, utilised in cases where the default OpenMP dialect handling cannot
-/// handle all cases for intermingled fir types and operations.
-void populateOpenMPFIRToLLVMConversionPatterns(
- LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns);
-
// declarative passes
#define GEN_PASS_REGISTRATION
#include "flang/Optimizer/CodeGen/CGPasses.h.inc"
diff --git a/flang/include/flang/Optimizer/CodeGen/CodeGenOpenMP.h b/flang/include/flang/Optimizer/CodeGen/CodeGenOpenMP.h
new file mode 100644
index 00000000000000..6c98fe1c3ffc62
--- /dev/null
+++ b/flang/include/flang/Optimizer/CodeGen/CodeGenOpenMP.h
@@ -0,0 +1,27 @@
+//=== Optimizer/CodeGen/CodeGenOpenMP.h - OpenMP code generation -*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef FORTRAN_OPTIMIZER_CODEGEN_CODEGENOPENMP_H
+#define FORTRAN_OPTIMIZER_CODEGEN_CODEGENOPENMP_H
+
+#include "mlir/Pass/Pass.h"
+#include "mlir/Pass/PassRegistry.h"
+
+namespace fir {
+class LLVMTypeConverter;
+
+/// Specialised conversion patterns of OpenMP operations for FIR to LLVM
+/// dialect, utilised in cases where the default OpenMP dialect handling cannot
+/// handle all cases for intermingled fir types and operations.
+void populateOpenMPFIRToLLVMConversionPatterns(
+ LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns);
+
+} // namespace fir
+
+
+#endif // FORTRAN_OPTIMIZER_CODEGEN_CODEGENOPENMP_H
\ No newline at end of file
diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index 9a45c88c828132..e04ac4b008d3f1 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -13,6 +13,7 @@
#include "flang/Optimizer/CodeGen/CodeGen.h"
#include "CGOps.h"
+#include "flang/Optimizer/CodeGen/CodeGenOpenMP.h"
#include "flang/Optimizer/Dialect/FIRAttr.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIRType.h"
diff --git a/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp b/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
index 3dadfe74817a45..b3ba695fa9be18 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGenOpenMP.cpp
@@ -1,7 +1,8 @@
-#include "flang/Optimizer/CodeGen/CodeGen.h"
+#include "flang/Optimizer/CodeGen/CodeGenOpenMP.h"
#include "flang/Optimizer/Builder/FIRBuilder.h"
#include "flang/Optimizer/Builder/LowLevelIntrinsics.h"
+#include "flang/Optimizer/CodeGen/CodeGen.h"
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#include "flang/Optimizer/Dialect/FIRType.h"
>From c9db5eec31c51fa3f3cbd081c744a4b04d4b2da2 Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Wed, 6 Dec 2023 08:46:47 -0600
Subject: [PATCH 11/19] Add minor adjustment of swapping to using BaseBoxType
checks
---
flang/lib/Lower/OpenMP.cpp | 4 ++--
flang/lib/Optimizer/Dialect/FIRType.cpp | 3 +--
2 files changed, 3 insertions(+), 4 deletions(-)
diff --git a/flang/lib/Lower/OpenMP.cpp b/flang/lib/Lower/OpenMP.cpp
index f6c78aec30edcd..fad90e853df606 100644
--- a/flang/lib/Lower/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP.cpp
@@ -1759,7 +1759,7 @@ static mlir::omp::MapInfoOp processDescriptorTypeMappings(
// !fir.ref<!fir.box<...>> to access the data we need to map we must
// perform an alloca and then store to it and retrieve the data from the new
// alloca.
- if (fir::isAssumedShape(fir::unwrapRefType(descriptorAddr.getType()))) {
+ if (mlir::isa<fir::BaseBoxType>(descriptorAddr.getType())) {
mlir::OpBuilder::InsertPoint insPt = firOpBuilder.saveInsertionPoint();
firOpBuilder.setInsertionPointToStart(firOpBuilder.getAllocaBlock());
descriptor =
@@ -1777,7 +1777,7 @@ static mlir::omp::MapInfoOp processDescriptorTypeMappings(
mapCaptureType),
mlir::omp::VariableCaptureKind::ByRef, descDataBaseAddr.getType()));
- // TODO: map the addendum segment of the descriptor, similarly to the abose
+ // TODO: map the addendum segment of the descriptor, similarly to the above
// base address/data pointer member.
return createMapInfoOp(
diff --git a/flang/lib/Optimizer/Dialect/FIRType.cpp b/flang/lib/Optimizer/Dialect/FIRType.cpp
index 9323a019dd7b55..ec15f807313321 100644
--- a/flang/lib/Optimizer/Dialect/FIRType.cpp
+++ b/flang/lib/Optimizer/Dialect/FIRType.cpp
@@ -331,8 +331,7 @@ bool isAllocatableOrPointerArray(mlir::Type ty) {
}
bool isTypeWithDescriptor(mlir::Type ty) {
- if (fir::isPointerType(ty) || fir::isAllocatableType(ty) ||
- fir::isAssumedShape(ty))
+ if (mlir::isa<fir::BaseBoxType>(unwrapRefType(ty)))
return true;
return false;
}
>From a3837fe9dd8418f8d7809482d443a085b7e792d3 Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Wed, 6 Dec 2023 09:29:17 -0600
Subject: [PATCH 12/19] Fix bad formatting
---
flang/include/flang/Optimizer/CodeGen/CodeGenOpenMP.h | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/flang/include/flang/Optimizer/CodeGen/CodeGenOpenMP.h b/flang/include/flang/Optimizer/CodeGen/CodeGenOpenMP.h
index 6c98fe1c3ffc62..eaef9ee130209d 100644
--- a/flang/include/flang/Optimizer/CodeGen/CodeGenOpenMP.h
+++ b/flang/include/flang/Optimizer/CodeGen/CodeGenOpenMP.h
@@ -1,4 +1,5 @@
-//=== Optimizer/CodeGen/CodeGenOpenMP.h - OpenMP code generation -*- C++ -*-===//
+//=== Optimizer/CodeGen/CodeGenOpenMP.h - OpenMP code generation -*- C++
+//-*-===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
@@ -23,5 +24,4 @@ void populateOpenMPFIRToLLVMConversionPatterns(
} // namespace fir
-
#endif // FORTRAN_OPTIMIZER_CODEGEN_CODEGENOPENMP_H
\ No newline at end of file
>From d3470221d1ccd86b5b14f1e5521a523ebc089917 Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Thu, 7 Dec 2023 07:23:48 -0600
Subject: [PATCH 13/19] Make sure we have a symbol before using it
---
flang/lib/Lower/OpenMP.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/flang/lib/Lower/OpenMP.cpp b/flang/lib/Lower/OpenMP.cpp
index fad90e853df606..83551fd6e4a3ca 100644
--- a/flang/lib/Lower/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP.cpp
@@ -1856,7 +1856,7 @@ bool ClauseProcessor::processMap(
auto origSymbol =
converter.getSymbolAddress(*getOmpObjectSymbol(ompObject));
mlir::Value mapOp, symAddr;
- if (fir::isTypeWithDescriptor(origSymbol.getType())) {
+ if (origSymbol && fir::isTypeWithDescriptor(origSymbol.getType())) {
symAddr = origSymbol;
mapOp = processDescriptorTypeMappings(
semanticsContext, stmtCtx, converter, clauseLocation,
>From 8e6329990e665af3bda4f4b1360ad7b564227c26 Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Wed, 3 Jan 2024 09:31:58 -0600
Subject: [PATCH 14/19] Fix rebase issues
---
flang/lib/Lower/OpenMP.cpp | 14 ++++++++------
1 file changed, 8 insertions(+), 6 deletions(-)
diff --git a/flang/lib/Lower/OpenMP.cpp b/flang/lib/Lower/OpenMP.cpp
index 83551fd6e4a3ca..6359dfae20f580 100644
--- a/flang/lib/Lower/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP.cpp
@@ -1848,11 +1848,13 @@ bool ClauseProcessor::processMap(
llvm::SmallVector<mlir::Value> bounds;
std::stringstream asFortran;
- Fortran::lower::AddrAndBoundsInfo info = Fortran::lower::gatherDataOperandAddrAndBounds<
- Fortran::parser::OmpObject, mlir::omp::DataBoundsType,
- mlir::omp::DataBoundsOp>(
- converter, firOpBuilder, semanticsContext, stmtCtx, ompObject,
- clauseLocation, asFortran, bounds, treatIndexAsSection);
+ Fortran::lower::AddrAndBoundsInfo info =
+ Fortran::lower::gatherDataOperandAddrAndBounds<
+ Fortran::parser::OmpObject, mlir::omp::DataBoundsOp,
+ mlir::omp::DataBoundsType>(
+ converter, firOpBuilder, semanticsContext, stmtCtx, ompObject,
+ clauseLocation, asFortran, bounds, treatIndexAsSection);
+
auto origSymbol =
converter.getSymbolAddress(*getOmpObjectSymbol(ompObject));
mlir::Value mapOp, symAddr;
@@ -2821,7 +2823,7 @@ genTargetOp(Fortran::lower::AbstractConverter &converter,
if (fir::isTypeWithDescriptor(baseOp.getType())) {
mapOp = processDescriptorTypeMappings(
semanticsContext, stmtCtx, converter, baseOp.getLoc(), baseOp,
- baseAddr, bounds, name.str(), mapFlag);
+ info.addr, bounds, name.str(), mapFlag);
} else {
mapOp = createMapInfoOp(
converter.getFirOpBuilder(), baseOp.getLoc(), baseOp,
>From bf2bdf3b390bdcfab1840b66a9d67f13d947ab00 Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Wed, 3 Jan 2024 16:04:00 -0600
Subject: [PATCH 15/19] [Flang][MLIR] Move majority of OpenMP lowering of
descriptor types into Opt Pass
---
.../flang/Optimizer/Transforms/Passes.h | 1 +
.../flang/Optimizer/Transforms/Passes.td | 11 ++
flang/include/flang/Tools/CLOptions.inc | 1 +
flang/lib/Lower/OpenMP.cpp | 129 +++++-------------
flang/lib/Optimizer/Transforms/CMakeLists.txt | 1 +
.../Transforms/OMPDescriptorMapInfoGen.cpp | 123 +++++++++++++++++
mlir/lib/Target/LLVMIR/ModuleTranslation.cpp | 2 +
7 files changed, 173 insertions(+), 95 deletions(-)
create mode 100644 flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp
diff --git a/flang/include/flang/Optimizer/Transforms/Passes.h b/flang/include/flang/Optimizer/Transforms/Passes.h
index 6970da8698ae84..aefb277f7966ba 100644
--- a/flang/include/flang/Optimizer/Transforms/Passes.h
+++ b/flang/include/flang/Optimizer/Transforms/Passes.h
@@ -76,6 +76,7 @@ std::unique_ptr<mlir::Pass>
createAlgebraicSimplificationPass(const mlir::GreedyRewriteConfig &config);
std::unique_ptr<mlir::Pass> createPolymorphicOpConversionPass();
+std::unique_ptr<mlir::Pass> createOMPDescriptorMapInfoGenPass();
std::unique_ptr<mlir::Pass> createOMPFunctionFilteringPass();
std::unique_ptr<mlir::OperationPass<mlir::ModuleOp>>
createOMPMarkDeclareTargetPass();
diff --git a/flang/include/flang/Optimizer/Transforms/Passes.td b/flang/include/flang/Optimizer/Transforms/Passes.td
index e3c45d41f04cc7..fe15d8bc53681b 100644
--- a/flang/include/flang/Optimizer/Transforms/Passes.td
+++ b/flang/include/flang/Optimizer/Transforms/Passes.td
@@ -318,6 +318,17 @@ def LoopVersioning : Pass<"loop-versioning", "mlir::func::FuncOp"> {
let dependentDialects = [ "fir::FIROpsDialect" ];
}
+def OMPDescriptorMapInfoGenPass
+ : Pass<"omp-descriptor-map-info-gen", "mlir::ModuleOp"> {
+ let summary = "expands OpenMP MapInfo operations containing descriptors";
+ let description = [{
+ Expands MapInfo operations containing descriptor types into multiple MapInfo's for each pointer element in
+ the descriptor that requires explicit individual mapping by the OpenMP runtime.
+ }];
+ let constructor = "::fir::createOMPDescriptorMapInfoGenPass()";
+ let dependentDialects = ["mlir::omp::OpenMPDialect"];
+}
+
def OMPMarkDeclareTargetPass
: Pass<"omp-mark-declare-target", "mlir::ModuleOp"> {
let summary = "Marks all functions called by an OpenMP declare target function as declare target";
diff --git a/flang/include/flang/Tools/CLOptions.inc b/flang/include/flang/Tools/CLOptions.inc
index 96d3869cd09391..8dee3074a5d4f8 100644
--- a/flang/include/flang/Tools/CLOptions.inc
+++ b/flang/include/flang/Tools/CLOptions.inc
@@ -274,6 +274,7 @@ inline void createHLFIRToFIRPassPipeline(
/// rather than the host device.
inline void createOpenMPFIRPassPipeline(
mlir::PassManager &pm, bool isTargetDevice) {
+ pm.addPass(fir::createOMPDescriptorMapInfoGenPass());
pm.addPass(fir::createOMPMarkDeclareTargetPass());
if (isTargetDevice)
pm.addPass(fir::createOMPFunctionFilteringPass());
diff --git a/flang/lib/Lower/OpenMP.cpp b/flang/lib/Lower/OpenMP.cpp
index 6359dfae20f580..e5178ac6a13102 100644
--- a/flang/lib/Lower/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP.cpp
@@ -1740,54 +1740,6 @@ createMapInfoOp(fir::FirOpBuilder &builder, mlir::Location loc,
return op;
}
-static mlir::omp::MapInfoOp processDescriptorTypeMappings(
- Fortran::semantics::SemanticsContext &semanticsContext,
- Fortran::lower::StatementContext &stmtCtx,
- Fortran::lower::AbstractConverter &converter, mlir::Location loc,
- mlir::Value descriptorAddr, mlir::Value descDataBaseAddr,
- mlir::ValueRange bounds, std::string asFortran,
- llvm::omp::OpenMPOffloadMappingFlags mapCaptureType) {
- llvm::SmallVector<mlir::Value> descriptorBaseAddrMembers;
- fir::FirOpBuilder &firOpBuilder = converter.getFirOpBuilder();
-
- mlir::Value descriptor = descriptorAddr;
-
- // The fir::BoxOffsetOp only works with !fir.ref<!fir.box<...>> types, as
- // allowing it to access non-reference box operations can cause some
- // problematic SSA IR. However, in the case of assumed shape's the type
- // is not a !fir.ref, in these cases to retrieve the appropriate
- // !fir.ref<!fir.box<...>> to access the data we need to map we must
- // perform an alloca and then store to it and retrieve the data from the new
- // alloca.
- if (mlir::isa<fir::BaseBoxType>(descriptorAddr.getType())) {
- mlir::OpBuilder::InsertPoint insPt = firOpBuilder.saveInsertionPoint();
- firOpBuilder.setInsertionPointToStart(firOpBuilder.getAllocaBlock());
- descriptor =
- firOpBuilder.create<fir::AllocaOp>(loc, descriptorAddr.getType());
- firOpBuilder.restoreInsertionPoint(insPt);
- firOpBuilder.create<fir::StoreOp>(loc, descriptorAddr, descriptor);
- }
-
- mlir::Value baseAddrAddr = firOpBuilder.create<fir::BoxOffsetOp>(
- loc, descriptor, fir::BoxFieldAttr::base_addr);
-
- descriptorBaseAddrMembers.push_back(createMapInfoOp(
- firOpBuilder, loc, descDataBaseAddr, baseAddrAddr, asFortran, bounds, {},
- static_cast<std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
- mapCaptureType),
- mlir::omp::VariableCaptureKind::ByRef, descDataBaseAddr.getType()));
-
- // TODO: map the addendum segment of the descriptor, similarly to the above
- // base address/data pointer member.
-
- return createMapInfoOp(
- firOpBuilder, loc, descriptor, mlir::Value{}, asFortran, {},
- descriptorBaseAddrMembers,
- static_cast<std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
- mapCaptureType),
- mlir::omp::VariableCaptureKind::ByRef, descriptor.getType());
-}
-
bool ClauseProcessor::processMap(
mlir::Location currentLocation, const llvm::omp::Directive &directive,
Fortran::semantics::SemanticsContext &semanticsContext,
@@ -1857,24 +1809,20 @@ bool ClauseProcessor::processMap(
auto origSymbol =
converter.getSymbolAddress(*getOmpObjectSymbol(ompObject));
- mlir::Value mapOp, symAddr;
- if (origSymbol && fir::isTypeWithDescriptor(origSymbol.getType())) {
+ mlir::Value symAddr = info.addr;
+ if (origSymbol && fir::isTypeWithDescriptor(origSymbol.getType()))
symAddr = origSymbol;
- mapOp = processDescriptorTypeMappings(
- semanticsContext, stmtCtx, converter, clauseLocation,
- origSymbol, info.addr, bounds, asFortran.str(), mapTypeBits);
- } else {
- // Explicit map captures are captured ByRef by default,
- // optimisation passes may alter this to ByCopy or other capture
- // types to optimise
- symAddr = info.addr;
- mapOp = createMapInfoOp(
- firOpBuilder, clauseLocation, info.addr, mlir::Value{},
- asFortran.str(), bounds, {},
- static_cast<std::underlying_type_t<
- llvm::omp::OpenMPOffloadMappingFlags>>(mapTypeBits),
- mlir::omp::VariableCaptureKind::ByRef, info.addr.getType());
- }
+
+ // Explicit map captures are captured ByRef by default,
+ // optimisation passes may alter this to ByCopy or other capture
+ // types to optimise
+ mlir::Value mapOp = createMapInfoOp(
+ firOpBuilder, clauseLocation, symAddr, mlir::Value{},
+ asFortran.str(), bounds, {},
+ static_cast<
+ std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
+ mapTypeBits),
+ mlir::omp::VariableCaptureKind::ByRef, symAddr.getType());
mapOperands.push_back(mapOp);
if (mapSymTypes)
@@ -1989,22 +1937,20 @@ bool ClauseProcessor::processMotionClauses(
auto origSymbol =
converter.getSymbolAddress(*getOmpObjectSymbol(ompObject));
- mlir::Value mapOp;
- if (origSymbol && fir::isTypeWithDescriptor(origSymbol.getType())) {
- mapOp = processDescriptorTypeMappings(
- semanticsContext, stmtCtx, converter, clauseLocation,
- origSymbol, info.addr, bounds, asFortran.str(), mapTypeBits);
- } else {
- // Explicit map captures are captured ByRef by default,
- // optimisation passes may alter this to ByCopy or other capture
- // types to optimise
- mapOp = createMapInfoOp(
- firOpBuilder, clauseLocation, info.addr, mlir::Value{},
- asFortran.str(), bounds, {},
- static_cast<std::underlying_type_t<
- llvm::omp::OpenMPOffloadMappingFlags>>(mapTypeBits),
- mlir::omp::VariableCaptureKind::ByRef, info.addr.getType());
- }
+ mlir::Value symAddr = info.addr;
+ if (origSymbol && fir::isTypeWithDescriptor(origSymbol.getType()))
+ symAddr = origSymbol;
+
+ // Explicit map captures are captured ByRef by default,
+ // optimisation passes may alter this to ByCopy or other capture
+ // types to optimise
+ mlir::Value mapOp = createMapInfoOp(
+ firOpBuilder, clauseLocation, symAddr, mlir::Value{},
+ asFortran.str(), bounds, {},
+ static_cast<
+ std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
+ mapTypeBits),
+ mlir::omp::VariableCaptureKind::ByRef, symAddr.getType());
mapOperands.push_back(mapOp);
}
@@ -2819,20 +2765,13 @@ genTargetOp(Fortran::lower::AbstractConverter &converter,
mapFlag |= llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_FROM;
}
- mlir::Value mapOp;
- if (fir::isTypeWithDescriptor(baseOp.getType())) {
- mapOp = processDescriptorTypeMappings(
- semanticsContext, stmtCtx, converter, baseOp.getLoc(), baseOp,
- info.addr, bounds, name.str(), mapFlag);
- } else {
- mapOp = createMapInfoOp(
- converter.getFirOpBuilder(), baseOp.getLoc(), baseOp,
- mlir::Value{}, name.str(), bounds, {},
- static_cast<
- std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
- mapFlag),
- captureKind, baseOp.getType());
- }
+ mlir::Value mapOp = createMapInfoOp(
+ converter.getFirOpBuilder(), baseOp.getLoc(), baseOp, mlir::Value{},
+ name.str(), bounds, {},
+ static_cast<
+ std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
+ mapFlag),
+ captureKind, baseOp.getType());
mapOperands.push_back(mapOp);
mapSymTypes.push_back(baseOp.getType());
diff --git a/flang/lib/Optimizer/Transforms/CMakeLists.txt b/flang/lib/Optimizer/Transforms/CMakeLists.txt
index fc067ad3585395..ba2e267996150e 100644
--- a/flang/lib/Optimizer/Transforms/CMakeLists.txt
+++ b/flang/lib/Optimizer/Transforms/CMakeLists.txt
@@ -17,6 +17,7 @@ add_flang_library(FIRTransforms
AddDebugFoundation.cpp
PolymorphicOpConversion.cpp
LoopVersioning.cpp
+ OMPDescriptorMapInfoGen.cpp
OMPFunctionFiltering.cpp
OMPMarkDeclareTarget.cpp
VScaleAttr.cpp
diff --git a/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp b/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp
new file mode 100644
index 00000000000000..30f05fb5f173a9
--- /dev/null
+++ b/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp
@@ -0,0 +1,123 @@
+#include "flang/Optimizer/Builder/FIRBuilder.h"
+#include "flang/Optimizer/Dialect/FIRType.h"
+#include "flang/Optimizer/Dialect/Support/KindMapping.h"
+#include "flang/Optimizer/Transforms/Passes.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
+#include "mlir/IR/BuiltinDialect.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/IR/SymbolTable.h"
+#include "mlir/Pass/Pass.h"
+#include "mlir/Support/LLVM.h"
+#include "llvm/ADT/SmallPtrSet.h"
+
+namespace fir {
+#define GEN_PASS_DEF_OMPDESCRIPTORMAPINFOGENPASS
+#include "flang/Optimizer/Transforms/Passes.h.inc"
+} // namespace fir
+
+namespace {
+class OMPDescriptorMapInfoGenPass
+ : public fir::impl::OMPDescriptorMapInfoGenPassBase<
+ OMPDescriptorMapInfoGenPass> {
+
+ mlir::omp::MapInfoOp
+ createMapInfoOp(fir::FirOpBuilder &builder, mlir::Location loc,
+ mlir::Value baseAddr, mlir::Value varPtrPtr, std::string name,
+ mlir::SmallVector<mlir::Value> bounds,
+ mlir::SmallVector<mlir::Value> members, uint64_t mapType,
+ mlir::omp::VariableCaptureKind mapCaptureType,
+ mlir::Type retTy, bool isVal = false) {
+ if (auto boxTy = baseAddr.getType().dyn_cast<fir::BaseBoxType>()) {
+ baseAddr = builder.create<fir::BoxAddrOp>(loc, baseAddr);
+ retTy = baseAddr.getType();
+ }
+
+ mlir::TypeAttr varType = mlir::TypeAttr::get(
+ llvm::cast<mlir::omp::PointerLikeType>(retTy).getElementType());
+
+ mlir::omp::MapInfoOp op = builder.create<mlir::omp::MapInfoOp>(
+ loc, retTy, baseAddr, varType, varPtrPtr, members, bounds,
+ builder.getIntegerAttr(builder.getIntegerType(64, false), mapType),
+ builder.getAttr<mlir::omp::VariableCaptureKindAttr>(mapCaptureType),
+ builder.getStringAttr(name));
+
+ return op;
+ }
+
+ void genDescriptorMemberMaps(mlir::omp::MapInfoOp op,
+ fir::FirOpBuilder &builder) {
+ llvm::SmallVector<mlir::Value> descriptorBaseAddrMembers;
+ mlir::Location loc = builder.getUnknownLoc();
+ mlir::Value descriptor = op.getVarPtr();
+
+ // If we enter this function, but the mapped type itself is not the
+ // descriptor, then it's likely the address of the descriptor so we
+ // must retrieve the descriptor SSA.
+ if (!fir::isTypeWithDescriptor(op.getVarType())) {
+ if (auto addrOp = mlir::dyn_cast_if_present<fir::BoxAddrOp>(
+ op.getVarPtr().getDefiningOp())) {
+ descriptor = addrOp.getVal();
+ }
+ }
+
+ // The fir::BoxOffsetOp only works with !fir.ref<!fir.box<...>> types, as
+ // allowing it to access non-reference box operations can cause some
+ // problematic SSA IR. However, in the case of assumed shape's the type
+ // is not a !fir.ref, in these cases to retrieve the appropriate
+ // !fir.ref<!fir.box<...>> to access the data we need to map we must
+ // perform an alloca and then store to it and retrieve the data from the new
+ // alloca.
+ if (mlir::isa<fir::BaseBoxType>(descriptor.getType())) {
+ mlir::OpBuilder::InsertPoint insPt = builder.saveInsertionPoint();
+ builder.setInsertionPointToStart(builder.getAllocaBlock());
+ auto alloca = builder.create<fir::AllocaOp>(loc, descriptor.getType());
+ builder.restoreInsertionPoint(insPt);
+ builder.create<fir::StoreOp>(loc, descriptor, alloca);
+ descriptor = alloca;
+ }
+
+ mlir::Value baseAddrAddr = builder.create<fir::BoxOffsetOp>(
+ loc, descriptor, fir::BoxFieldAttr::base_addr);
+
+ descriptorBaseAddrMembers.push_back(createMapInfoOp(
+ builder, loc, baseAddrAddr, {}, "", op.getBounds(), {},
+ op.getMapType().value(), mlir::omp::VariableCaptureKind::ByRef,
+ fir::unwrapRefType(baseAddrAddr.getType())));
+
+ // TODO: map the addendum segment of the descriptor, similarly to the above
+ // base address/data pointer member.
+
+ op.getVarPtrMutable().assign(descriptor);
+ op.setVarType(fir::unwrapRefType(descriptor.getType()));
+ op.getMembersMutable().assign(descriptorBaseAddrMembers);
+ op.getBoundsMutable().assign(llvm::SmallVector<mlir::Value>{});
+ }
+
+ // This pass executes on mlir::ModuleOp's finding omp::MapInfoOp's containing
+ // descriptor based types (allocatables, pointers, assumed shape etc.) and
+ // expanding them into multiple omp::MapInfoOp's for each pointer member
+ // contained within the descriptor.
+ void runOnOperation() override {
+ fir::KindMapping kindMap = fir::getKindMapping(getOperation());
+ fir::FirOpBuilder builder{getOperation(), std::move(kindMap)};
+
+ getOperation()->walk([&](mlir::omp::MapInfoOp op) {
+ if (fir::isTypeWithDescriptor(op.getVarType()) ||
+ mlir::isa<fir::BoxAddrOp>(op.getVarPtr().getDefiningOp())) {
+ builder.setInsertionPoint(op);
+ genDescriptorMemberMaps(op, builder);
+ }
+ });
+ }
+};
+
+} // namespace
+
+namespace fir {
+std::unique_ptr<mlir::Pass> createOMPDescriptorMapInfoGenPass() {
+ return std::make_unique<OMPDescriptorMapInfoGenPass>();
+}
+} // namespace fir
diff --git a/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp b/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
index 1722d74c08b628..d6c6c7f86f85c4 100644
--- a/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
@@ -1584,6 +1584,8 @@ mlir::translateModuleToLLVMIR(Operation *module, llvm::LLVMContext &llvmContext,
if (failed(translator.convertFunctions()))
return nullptr;
+ // translator.llvmModule->dump();
+
if (llvm::verifyModule(*translator.llvmModule, &llvm::errs()))
return nullptr;
>From 3e4d168e747627a95d73df1f572f0de54fd71e97 Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Thu, 4 Jan 2024 12:11:30 -0600
Subject: [PATCH 16/19] [Flang][MLIR] Insert descriptor implicit members into
map operands and BlockArgs
---
.../Transforms/OMPDescriptorMapInfoGen.cpp | 38 ++++++-
mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td | 19 ++--
.../Dialect/OpenMP/OpenMPOpsInterfaces.td | 24 ++++-
mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp | 2 +-
.../OpenMP/OpenMPToLLVMIRTranslation.cpp | 101 +++++++-----------
5 files changed, 110 insertions(+), 74 deletions(-)
diff --git a/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp b/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp
index 30f05fb5f173a9..a5f5500f58b624 100644
--- a/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp
+++ b/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp
@@ -12,6 +12,7 @@
#include "mlir/Pass/Pass.h"
#include "mlir/Support/LLVM.h"
#include "llvm/ADT/SmallPtrSet.h"
+#include <iterator>
namespace fir {
#define GEN_PASS_DEF_OMPDESCRIPTORMAPINFOGENPASS
@@ -48,7 +49,8 @@ class OMPDescriptorMapInfoGenPass
}
void genDescriptorMemberMaps(mlir::omp::MapInfoOp op,
- fir::FirOpBuilder &builder) {
+ fir::FirOpBuilder &builder,
+ mlir::Operation *target) {
llvm::SmallVector<mlir::Value> descriptorBaseAddrMembers;
mlir::Location loc = builder.getUnknownLoc();
mlir::Value descriptor = op.getVarPtr();
@@ -92,8 +94,36 @@ class OMPDescriptorMapInfoGenPass
op.getVarPtrMutable().assign(descriptor);
op.setVarType(fir::unwrapRefType(descriptor.getType()));
- op.getMembersMutable().assign(descriptorBaseAddrMembers);
+ op.getMembersMutable().append(descriptorBaseAddrMembers);
op.getBoundsMutable().assign(llvm::SmallVector<mlir::Value>{});
+
+ // could use a template to generalise to other TargetOps
+ if (auto mapClauseOwner =
+ llvm::dyn_cast<mlir::omp::MapClauseOwningOpInterface>(target)) {
+ llvm::SmallVector<mlir::Value> newMapOps;
+ for (size_t i = 0; i < mapClauseOwner.getMapOperands().size(); ++i) {
+ if (mapClauseOwner.getMapOperands()[i] == op) {
+ // Push new implicit maps generated for the descriptor.
+ newMapOps.push_back(descriptorBaseAddrMembers[0]);
+
+ // for TargetOp's which have IsolatedFromAbove we must align the
+ // new additional map operand with an appropriate BlockArgument,
+ // as the printing and later processing currently requires a 1:1
+ // mapping of BlockArgs to MapInfoOp's at the same placement in
+ // each array (BlockArgs and MapOperands).
+ if (auto targetOp = llvm::dyn_cast<mlir::omp::TargetOp>(target)) {
+ targetOp.getRegion().insertArgument(
+ i, descriptorBaseAddrMembers[0].getType(), loc);
+ }
+
+ newMapOps.push_back(mapClauseOwner.getMapOperands()[i]);
+ } else {
+ newMapOps.push_back(mapClauseOwner.getMapOperands()[i]);
+ }
+ }
+
+ mapClauseOwner.getMapOperandsMutable().assign(newMapOps);
+ }
}
// This pass executes on mlir::ModuleOp's finding omp::MapInfoOp's containing
@@ -108,7 +138,9 @@ class OMPDescriptorMapInfoGenPass
if (fir::isTypeWithDescriptor(op.getVarType()) ||
mlir::isa<fir::BoxAddrOp>(op.getVarPtr().getDefiningOp())) {
builder.setInsertionPoint(op);
- genDescriptorMemberMaps(op, builder);
+ // Currently a MapInfoOp argument can only show up on a single target
+ // user so we can retrieve and use the first user.
+ genDescriptorMemberMaps(op, builder, *op->getUsers().begin());
}
});
}
diff --git a/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td b/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
index 89096650d84142..8f1b162edfcc85 100644
--- a/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
+++ b/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
@@ -1232,7 +1232,8 @@ def MapInfoOp : OpenMP_Op<"map_info", [AttrSizedOperandSegments]> {
// 2.14.2 target data Construct
//===---------------------------------------------------------------------===//
-def Target_DataOp: OpenMP_Op<"target_data", [AttrSizedOperandSegments]>{
+def Target_DataOp: OpenMP_Op<"target_data", [AttrSizedOperandSegments,
+ MapClauseOwningOpInterface]>{
let summary = "target data construct";
let description = [{
Map variables to a device data environment for the extent of the region.
@@ -1289,7 +1290,8 @@ def Target_DataOp: OpenMP_Op<"target_data", [AttrSizedOperandSegments]>{
//===---------------------------------------------------------------------===//
def Target_EnterDataOp: OpenMP_Op<"target_enter_data",
- [AttrSizedOperandSegments]>{
+ [AttrSizedOperandSegments,
+ MapClauseOwningOpInterface]>{
let summary = "target enter data construct";
let description = [{
The target enter data directive specifies that variables are mapped to
@@ -1335,7 +1337,8 @@ def Target_EnterDataOp: OpenMP_Op<"target_enter_data",
//===---------------------------------------------------------------------===//
def Target_ExitDataOp: OpenMP_Op<"target_exit_data",
- [AttrSizedOperandSegments]>{
+ [AttrSizedOperandSegments,
+ MapClauseOwningOpInterface]>{
let summary = "target exit data construct";
let description = [{
The target exit data directive specifies that variables are mapped to a
@@ -1381,7 +1384,8 @@ def Target_ExitDataOp: OpenMP_Op<"target_exit_data",
//===---------------------------------------------------------------------===//
def Target_UpdateDataOp: OpenMP_Op<"target_update_data",
- [AttrSizedOperandSegments]>{
+ [AttrSizedOperandSegments,
+ MapClauseOwningOpInterface]>{
let summary = "target update data construct";
let description = [{
The target update directive makes the corresponding list items in the device
@@ -1413,13 +1417,13 @@ def Target_UpdateDataOp: OpenMP_Op<"target_update_data",
let arguments = (ins Optional<I1>:$if_expr,
Optional<AnyInteger>:$device,
UnitAttr:$nowait,
- Variadic<OpenMP_PointerLikeType>:$motion_operands);
+ Variadic<OpenMP_PointerLikeType>:$map_operands);
let assemblyFormat = [{
oilist(`if` `(` $if_expr `:` type($if_expr) `)`
| `device` `(` $device `:` type($device) `)`
| `nowait` $nowait
- | `motion_entries` `(` $motion_operands `:` type($motion_operands) `)`
+ | `motion_entries` `(` $map_operands `:` type($map_operands) `)`
) attr-dict
}];
@@ -1430,7 +1434,8 @@ def Target_UpdateDataOp: OpenMP_Op<"target_update_data",
// 2.14.5 target construct
//===----------------------------------------------------------------------===//
-def TargetOp : OpenMP_Op<"target",[IsolatedFromAbove, OutlineableOpenMPOpInterface, AttrSizedOperandSegments]> {
+def TargetOp : OpenMP_Op<"target",[IsolatedFromAbove, MapClauseOwningOpInterface,
+ OutlineableOpenMPOpInterface, AttrSizedOperandSegments]> {
let summary = "target construct";
let description = [{
The target construct includes a region of code which is to be executed
diff --git a/mlir/include/mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td b/mlir/include/mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td
index 77001fc816cf91..1567143e873c50 100644
--- a/mlir/include/mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td
+++ b/mlir/include/mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td
@@ -18,7 +18,7 @@ include "mlir/IR/OpBase.td"
def OutlineableOpenMPOpInterface : OpInterface<"OutlineableOpenMPOpInterface"> {
let description = [{
OpenMP operations whose region will be outlined will implement this
- interface. These operations will
+ interface.
}];
let cppNamespace = "::mlir::omp";
@@ -31,6 +31,28 @@ def OutlineableOpenMPOpInterface : OpInterface<"OutlineableOpenMPOpInterface"> {
];
}
+def MapClauseOwningOpInterface : OpInterface<"MapClauseOwningOpInterface"> {
+ let description = [{
+ OpenMP operations which own a list of omp::MapInfoOp's implement this interface
+ to allow generic access to deal with map operands to more easily manipulate
+ this class of operations.
+ }];
+
+ let cppNamespace = "::mlir::omp";
+
+ let methods = [
+ InterfaceMethod<"Get map operands", "::mlir::OperandRange", "getMapOperands",
+ (ins), [{
+ return $_op.getMapOperands();
+ }]>,
+ InterfaceMethod<"Get mutable map operands", "::mlir::MutableOperandRange",
+ "getMapOperandsMutable",
+ (ins), [{
+ return $_op.getMapOperandsMutable();
+ }]>,
+ ];
+}
+
def ReductionClauseInterface : OpInterface<"ReductionClauseInterface"> {
let description = [{
OpenMP operations that support reduction clause have this interface.
diff --git a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
index 6e69cd0d386bd2..d9611f21cd8d63 100644
--- a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
+++ b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
@@ -980,7 +980,7 @@ LogicalResult ExitDataOp::verify() {
}
LogicalResult UpdateDataOp::verify() {
- return verifyMapClause(*this, getMotionOperands());
+ return verifyMapClause(*this, getMapOperands());
}
LogicalResult TargetOp::verify() {
diff --git a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
index aea1a1bfedbcba..e8982f36a4e85c 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
@@ -1640,6 +1640,7 @@ getRefPtrIfDeclareTarget(mlir::Value value,
// value) more than neccessary.
struct MapInfoData : llvm::OpenMPIRBuilder::MapInfosTy {
llvm::SmallVector<bool, 4> IsDeclareTarget;
+ llvm::SmallVector<bool, 4> IsAMember;
llvm::SmallVector<mlir::Operation *, 4> MapClause;
llvm::SmallVector<llvm::Value *, 4> OriginalValue;
// Stripped off array/pointer to get the underlying
@@ -1728,13 +1729,11 @@ void collectMapDataFromMapOperands(MapInfoData &mapData,
LLVM::ModuleTranslation &moduleTranslation,
DataLayout &dl,
llvm::IRBuilderBase &builder) {
- auto mapFill = [&](mlir::Value mapValue) {
- assert(mlir::isa<mlir::omp::MapInfoOp>(mapValue.getDefiningOp()) &&
- "missing map info operation or incorrect map info operation type");
+ for (mlir::Value mapValue : mapOperands) {
if (auto mapOp = mlir::dyn_cast_if_present<mlir::omp::MapInfoOp>(
mapValue.getDefiningOp())) {
- mapData.OriginalValue.push_back(moduleTranslation.lookupValue(
- mapOp.getVarPtrPtr() ? mapOp.getVarPtrPtr() : mapOp.getVarPtr()));
+ mapData.OriginalValue.push_back(
+ moduleTranslation.lookupValue(mapOp.getVarPtr()));
mapData.Pointers.push_back(mapData.OriginalValue.back());
if (llvm::Value *refPtr =
@@ -1759,41 +1758,27 @@ void collectMapDataFromMapOperands(MapInfoData &mapData,
mapOp.getLoc(), *moduleTranslation.getOpenMPBuilder()));
mapData.DevicePointers.push_back(
llvm::OpenMPIRBuilder::DeviceInfoTy::None);
- }
- };
- // In the case of Fortran descriptors some members get added implicitly
- // after the target region has been generated during CodeGen lowering
- // which prevents them from being added trivially to the target region
- // as map arguments, we must handle this case here by generating
- // MapInfoData for them.
- // TODO: Revisit this when implementing derived types explicit member
- // mapping, we likely want to represent these identically to simplify
- // the overall lowering
- SmallVector<Value> mapMemberOperands;
- for (size_t i = 0; i < mapOperands.size(); ++i) {
- auto mapInfoOp =
- mlir::dyn_cast<mlir::omp::MapInfoOp>(mapOperands[i].getDefiningOp());
- for (auto members : mapInfoOp.getMembers()) {
- if (!std::any_of(mapOperands.begin(), mapOperands.end(),
- [&](auto mapOp) {
- return mapOp.getDefiningOp() ==
- members.getDefiningOp();
- }) &&
- !std::any_of(mapMemberOperands.begin(), mapMemberOperands.end(),
- [&](auto mapOp) {
- return mapOp.getDefiningOp() ==
- members.getDefiningOp();
- }))
- mapMemberOperands.push_back(members);
+ // Check if this is a member mapping and correctly assign that it is, if
+ // it is a member of a larger object.
+ // TODO: Need better handling of members, and distinguishing of members
+ // that are implicitly allocated on device vs explicitly passed in as
+ // arguments.
+ // TODO: May require some further additions to support nested record
+ // types, i.e. member maps that can have member maps.
+ mapData.IsAMember.push_back(false);
+ for (mlir::Value mapValue : mapOperands) {
+ if (auto map = mlir::dyn_cast_if_present<mlir::omp::MapInfoOp>(
+ mapValue.getDefiningOp())) {
+ for (auto member : map.getMembers()) {
+ if (member == mapOp) {
+ mapData.IsAMember.back() = true;
+ }
+ }
+ }
+ }
}
}
-
- for (mlir::Value mapValue : mapOperands)
- mapFill(mapValue);
-
- for (mlir::Value mapValue : mapMemberOperands)
- mapFill(mapValue);
}
static void processMapWithMembersOf(
@@ -1960,35 +1945,22 @@ static void genMapInfos(llvm::IRBuilderBase &builder,
combinedInfo.Names.clear();
};
- llvm::SmallVector<size_t, 4> primaryMapIdx;
- for (size_t i = 0; i < mapData.MapClause.size(); ++i) {
- primaryMapIdx.push_back(i);
- }
-
- // TODO: Handle nested MembersOf, currently only cares about the first level
- // of nesting (all that was relevant for Fortran descriptors), but a slight
- // refactoring of mapInfoData to hold nestings or membersOf may be a better
- // approach to simplify things.
- for (size_t i = 0; i < mapData.MapClause.size(); ++i) {
- auto mapInfoOp = mlir::dyn_cast<mlir::omp::MapInfoOp>(mapData.MapClause[i]);
- for (auto member : mapInfoOp.getMembers()) {
- for (size_t j = 0; j < primaryMapIdx.size(); j++) {
- if (member.getDefiningOp() == mapData.MapClause[primaryMapIdx[j]]) {
- primaryMapIdx.erase(&primaryMapIdx[j]);
- j--;
- }
- }
- }
- }
-
// We operate under the assumption that all vectors that are
// required in MapInfoData are of equal lengths (either filled with
// default constructed data or appropiate information) so we can
// utilise the size from any component of MapInfoData, if we can't
// something is missing from the initial MapInfoData construction.
- for (unsigned long i : primaryMapIdx) {
- auto mapInfoOp = mlir::dyn_cast<mlir::omp::MapInfoOp>(mapData.MapClause[i]);
+ for (size_t i = 0; i < mapData.MapClause.size(); ++i) {
+ // NOTE/TODO: We currently do not handle member mapping seperately from it's
+ // parent or explicit mapping of a parent and member in the same operation,
+ // this will need to change in the near future, for now we primarily handle
+ // descriptor mapping from fortran, generalised as mapping record types
+ // with implicit member maps. This lowering needs further generalisation to
+ // fully support fortran derived types, and C/C++ structures and classes.
+ if (mapData.IsAMember[i])
+ continue;
+ auto mapInfoOp = mlir::dyn_cast<mlir::omp::MapInfoOp>(mapData.MapClause[i]);
if (!mapInfoOp.getMembers().empty()) {
processMapWithMembersOf(moduleTranslation, builder, *ompBuilder, dl,
combinedInfo, mapData, i, isTargetParams);
@@ -2136,7 +2108,7 @@ convertOmpTargetData(Operation *op, llvm::IRBuilderBase &builder,
deviceID = intAttr.getInt();
RTLFn = llvm::omp::OMPRTL___tgt_target_data_update_mapper;
- mapOperands = updateDataOp.getMotionOperands();
+ mapOperands = updateDataOp.getMapOperands();
return success();
})
.Default([&](Operation *op) {
@@ -2633,7 +2605,12 @@ convertOmpTarget(Operation &opInst, llvm::IRBuilderBase &builder,
llvm::SmallVector<llvm::Value *, 4> kernelInput;
for (size_t i = 0; i < mapOperands.size(); ++i) {
// declare target arguments are not passed to kernels as arguments
- if (!mapData.IsDeclareTarget[i])
+ // TODO: We currently do not handle cases where a member is explicitly
+ // passed in as an argument, this will likley need to be handled in
+ // the near future, rather than using IsAMember, it may be better to
+ // test if the relevant BlockArg is used within the target region and
+ // then use that as a basis for exclusion in the kernel inputs.
+ if (!mapData.IsDeclareTarget[i] && !mapData.IsAMember[i])
kernelInput.push_back(mapData.OriginalValue[i]);
}
>From f917a255755fe53872079a7ad2122a22cfef222a Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Thu, 4 Jan 2024 18:55:51 -0600
Subject: [PATCH 17/19] [Flang][MLIR] Update tests based on moving to Opt pass
method and remove createMapInfoOp use from opt pass
---
.../Transforms/OMPDescriptorMapInfoGen.cpp | 51 +++++++------------
flang/test/Lower/OpenMP/FIR/array-bounds.f90 | 7 ++-
flang/test/Lower/OpenMP/FIR/target.f90 | 6 +--
.../Lower/OpenMP/allocatable-array-bounds.f90 | 24 ++++-----
flang/test/Lower/OpenMP/allocatable-map.f90 | 8 ++-
flang/test/Lower/OpenMP/array-bounds.f90 | 9 ++--
flang/test/Lower/OpenMP/target.f90 | 6 +--
...target-fortran-allocatable-types-host.mlir | 10 ++--
8 files changed, 50 insertions(+), 71 deletions(-)
diff --git a/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp b/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp
index a5f5500f58b624..17e0ce05143e03 100644
--- a/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp
+++ b/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp
@@ -24,30 +24,6 @@ class OMPDescriptorMapInfoGenPass
: public fir::impl::OMPDescriptorMapInfoGenPassBase<
OMPDescriptorMapInfoGenPass> {
- mlir::omp::MapInfoOp
- createMapInfoOp(fir::FirOpBuilder &builder, mlir::Location loc,
- mlir::Value baseAddr, mlir::Value varPtrPtr, std::string name,
- mlir::SmallVector<mlir::Value> bounds,
- mlir::SmallVector<mlir::Value> members, uint64_t mapType,
- mlir::omp::VariableCaptureKind mapCaptureType,
- mlir::Type retTy, bool isVal = false) {
- if (auto boxTy = baseAddr.getType().dyn_cast<fir::BaseBoxType>()) {
- baseAddr = builder.create<fir::BoxAddrOp>(loc, baseAddr);
- retTy = baseAddr.getType();
- }
-
- mlir::TypeAttr varType = mlir::TypeAttr::get(
- llvm::cast<mlir::omp::PointerLikeType>(retTy).getElementType());
-
- mlir::omp::MapInfoOp op = builder.create<mlir::omp::MapInfoOp>(
- loc, retTy, baseAddr, varType, varPtrPtr, members, bounds,
- builder.getIntegerAttr(builder.getIntegerType(64, false), mapType),
- builder.getAttr<mlir::omp::VariableCaptureKindAttr>(mapCaptureType),
- builder.getStringAttr(name));
-
- return op;
- }
-
void genDescriptorMemberMaps(mlir::omp::MapInfoOp op,
fir::FirOpBuilder &builder,
mlir::Operation *target) {
@@ -84,13 +60,20 @@ class OMPDescriptorMapInfoGenPass
mlir::Value baseAddrAddr = builder.create<fir::BoxOffsetOp>(
loc, descriptor, fir::BoxFieldAttr::base_addr);
- descriptorBaseAddrMembers.push_back(createMapInfoOp(
- builder, loc, baseAddrAddr, {}, "", op.getBounds(), {},
- op.getMapType().value(), mlir::omp::VariableCaptureKind::ByRef,
- fir::unwrapRefType(baseAddrAddr.getType())));
-
- // TODO: map the addendum segment of the descriptor, similarly to the above
- // base address/data pointer member.
+ descriptorBaseAddrMembers.push_back(builder.create<mlir::omp::MapInfoOp>(
+ loc, baseAddrAddr.getType(), baseAddrAddr,
+ llvm::cast<mlir::omp::PointerLikeType>(
+ fir::unwrapRefType(baseAddrAddr.getType()))
+ .getElementType(),
+ mlir::Value{}, mlir::SmallVector<mlir::Value>{}, op.getBounds(),
+ builder.getIntegerAttr(builder.getIntegerType(64, false),
+ op.getMapType().value()),
+ builder.getAttr<mlir::omp::VariableCaptureKindAttr>(
+ mlir::omp::VariableCaptureKind::ByRef),
+ builder.getStringAttr("")));
+
+ // TODO: map the addendum segment of the descriptor, similarly to the
+ // above base address/data pointer member.
op.getVarPtrMutable().assign(descriptor);
op.setVarType(fir::unwrapRefType(descriptor.getType()));
@@ -134,9 +117,13 @@ class OMPDescriptorMapInfoGenPass
fir::KindMapping kindMap = fir::getKindMapping(getOperation());
fir::FirOpBuilder builder{getOperation(), std::move(kindMap)};
+ // 2) look into adding some documentation on this.
+ // 3) fix tests.
+
getOperation()->walk([&](mlir::omp::MapInfoOp op) {
if (fir::isTypeWithDescriptor(op.getVarType()) ||
- mlir::isa<fir::BoxAddrOp>(op.getVarPtr().getDefiningOp())) {
+ mlir::isa_and_present<fir::BoxAddrOp>(
+ op.getVarPtr().getDefiningOp())) {
builder.setInsertionPoint(op);
// Currently a MapInfoOp argument can only show up on a single target
// user so we can retrieve and use the first user.
diff --git a/flang/test/Lower/OpenMP/FIR/array-bounds.f90 b/flang/test/Lower/OpenMP/FIR/array-bounds.f90
index ec0ecd3c90bb9e..e43dc15841ec5e 100644
--- a/flang/test/Lower/OpenMP/FIR/array-bounds.f90
+++ b/flang/test/Lower/OpenMP/FIR/array-bounds.f90
@@ -46,11 +46,10 @@ module assumed_array_routines
!ALL: %[[DIMS1:.*]]:3 = fir.box_dims %arg0, %[[C0_1]] : (!fir.box<!fir.array<?xi32>>, index) -> (index, index, index)
!ALL: %[[BOUNDS:.*]] = omp.bounds lower_bound(%[[C3]] : index) upper_bound(%[[C4]] : index) extent(%[[DIMS1]]#1 : index) stride(%[[DIMS0]]#2 : index) start_idx(%[[C0]] : index) {stride_in_bytes = true}
!ALL: %[[BOXADDRADDR:.*]] = fir.box_offset %0 base_addr : (!fir.ref<!fir.box<!fir.array<?xi32>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
-!ALL: %[[BOXADDR:.*]] = fir.box_addr %[[ARG0]] : (!fir.box<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
-!ALL: %[[MAP_MEMBER:.*]] = omp.map_info var_ptr(%[[BOXADDR]] : !fir.ref<!fir.array<?xi32>>, !fir.array<?xi32>) var_ptr_ptr(%[[BOXADDRADDR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS]]) -> !fir.ref<!fir.array<?xi32>> {name = "arr_read_write(2:5)"}
-!ALL: %[[MAP:.*]] = omp.map_info var_ptr(%0 : !fir.ref<!fir.box<!fir.array<?xi32>>>, !fir.box<!fir.array<?xi32>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_MEMBER]] : !fir.ref<!fir.array<?xi32>>) -> !fir.ref<!fir.box<!fir.array<?xi32>>> {name = "arr_read_write(2:5)"}
+!ALL: %[[MAP_MEMBER:.*]] = omp.map_info var_ptr(%[[BOXADDRADDR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>, !fir.array<?xi32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS]]) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>> {name = ""}
+!ALL: %[[MAP:.*]] = omp.map_info var_ptr(%0 : !fir.ref<!fir.box<!fir.array<?xi32>>>, !fir.box<!fir.array<?xi32>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_MEMBER]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) -> !fir.ref<!fir.array<?xi32>> {name = "arr_read_write(2:5)"}
!ALL: %[[MAP2:.*]] = omp.map_info var_ptr(%[[ALLOCA]] : !fir.ref<i32>, i32) map_clauses(implicit, exit_release_or_enter_alloc) capture(ByCopy) -> !fir.ref<i32> {name = "i"}
-!ALL: omp.target map_entries(%[[MAP]] -> %{{.*}}, %[[MAP2]] -> %{{.*}} : !fir.ref<!fir.box<!fir.array<?xi32>>>, !fir.ref<i32>) {
+!ALL: omp.target map_entries(%[[MAP_MEMBER]] -> %{{.*}}, %[[MAP]] -> %{{.*}}, %[[MAP2]] -> %{{.*}} : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>, !fir.ref<!fir.array<?xi32>>, !fir.ref<i32>) {
subroutine assumed_shape_array(arr_read_write)
integer, intent(inout) :: arr_read_write(:)
diff --git a/flang/test/Lower/OpenMP/FIR/target.f90 b/flang/test/Lower/OpenMP/FIR/target.f90
index 2018958b4e28cf..2d039bcea29de6 100644
--- a/flang/test/Lower/OpenMP/FIR/target.f90
+++ b/flang/test/Lower/OpenMP/FIR/target.f90
@@ -450,9 +450,9 @@ end subroutine omp_target_device_ptr
subroutine omp_target_device_addr
integer, pointer :: a
!CHECK: %[[VAL_0:.*]] = fir.alloca !fir.box<!fir.ptr<i32>> {bindc_name = "a", uniq_name = "_QFomp_target_device_addrEa"}
- !CHECK: %[[MAP_MEMBERS:.*]] = omp.map_info var_ptr({{.*}} : !fir.ptr<i32>, i32) var_ptr_ptr({{.*}} : !fir.llvm_ptr<!fir.ref<i32>>) map_clauses(tofrom) capture(ByRef) -> !fir.ptr<i32> {name = "a"}
- !CHECK: %[[MAP:.*]] = omp.map_info var_ptr({{.*}} : !fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.box<!fir.ptr<i32>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_MEMBERS]] : !fir.ptr<i32>) -> !fir.ref<!fir.box<!fir.ptr<i32>>> {name = "a"}
- !CHECK: omp.target_data map_entries(%[[MAP]] : {{.*}}) use_device_addr(%[[VAL_0]] : !fir.ref<!fir.box<!fir.ptr<i32>>>) {
+ !CHECK: %[[MAP_MEMBERS:.*]] = omp.map_info var_ptr({{.*}} : !fir.llvm_ptr<!fir.ref<i32>>, i32) map_clauses(tofrom) capture(ByRef) -> !fir.llvm_ptr<!fir.ref<i32>> {name = ""}
+ !CHECK: %[[MAP:.*]] = omp.map_info var_ptr({{.*}} : !fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.box<!fir.ptr<i32>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_MEMBERS]] : !fir.llvm_ptr<!fir.ref<i32>>) -> !fir.ref<!fir.box<!fir.ptr<i32>>> {name = "a"}
+ !CHECK: omp.target_data map_entries(%[[MAP_MEMBERS]], %[[MAP]] : {{.*}}) use_device_addr(%[[VAL_0]] : !fir.ref<!fir.box<!fir.ptr<i32>>>) {
!$omp target data map(tofrom: a) use_device_addr(a)
!CHECK: ^bb0(%[[VAL_1:.*]]: !fir.ref<!fir.box<!fir.ptr<i32>>>):
!CHECK: {{.*}} = fir.load %[[VAL_1]] : !fir.ref<!fir.box<!fir.ptr<i32>>>
diff --git a/flang/test/Lower/OpenMP/allocatable-array-bounds.f90 b/flang/test/Lower/OpenMP/allocatable-array-bounds.f90
index 5bf84dc3d5962a..fa50c63b40099a 100644
--- a/flang/test/Lower/OpenMP/allocatable-array-bounds.f90
+++ b/flang/test/Lower/OpenMP/allocatable-array-bounds.f90
@@ -22,10 +22,9 @@
!HOST: %[[CONSTANT_3:.*]] = arith.constant 0 : index
!HOST: %[[BOX_3:.*]]:3 = fir.box_dims %[[LOAD_3]], %[[CONSTANT_3]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>, index) -> (index, index, index)
!HOST: %[[BOUNDS_1:.*]] = omp.bounds lower_bound(%[[LB_1]] : index) upper_bound(%[[UB_1]] : index) extent(%[[BOX_3]]#1 : index) stride(%[[BOX_2]]#2 : index) start_idx(%[[BOX_1]]#0 : index) {stride_in_bytes = true}
-!HOST: %[[VAR_PTR_PTR:.*]] = fir.box_offset %[[DECLARE_1]]#1 base_addr : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
-!HOST: %[[VAR_PTR:.*]] = fir.box_addr %[[LOAD_1]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>) -> !fir.heap<!fir.array<?xi32>>
-!HOST: %[[MAP_INFO_MEMBER:.*]] = omp.map_info var_ptr(%[[VAR_PTR]] : !fir.heap<!fir.array<?xi32>>, !fir.array<?xi32>) var_ptr_ptr(%[[VAR_PTR_PTR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS_1]]) -> !fir.heap<!fir.array<?xi32>> {name = "sp_read(2:5)"}
-!HOST: %[[MAP_INFO_1:.*]] = omp.map_info var_ptr(%[[DECLARE_1]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, !fir.box<!fir.heap<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_INFO_MEMBER]] : !fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> {name = "sp_read(2:5)"}
+!HOST: %[[VAR_PTR:.*]] = fir.box_offset %[[DECLARE_1]]#1 base_addr : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
+!HOST: %[[MAP_INFO_MEMBER:.*]] = omp.map_info var_ptr(%[[VAR_PTR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>, !fir.array<?xi32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS_1]]) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>> {name = ""}
+!HOST: %[[MAP_INFO_1:.*]] = omp.map_info var_ptr(%[[DECLARE_1]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, !fir.box<!fir.heap<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_INFO_MEMBER]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> {name = "sp_read(2:5)"}
!HOST: %[[LOAD_3:.*]] = fir.load %[[DECLARE_2]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
!HOST: %[[LOAD_4:.*]] = fir.load %[[DECLARE_2]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
@@ -41,10 +40,9 @@
!HOST: %[[CONSTANT_5:.*]] = arith.constant 0 : index
!HOST: %[[BOX_5:.*]]:3 = fir.box_dims %[[LOAD_5]], %[[CONSTANT_5]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>, index) -> (index, index, index)
!HOST: %[[BOUNDS_2:.*]] = omp.bounds lower_bound(%[[LB_2]] : index) upper_bound(%[[UB_2]] : index) extent(%[[BOX_5]]#1 : index) stride(%[[BOX_4]]#2 : index) start_idx(%[[BOX_3]]#0 : index) {stride_in_bytes = true}
-!HOST: %[[VAR_PTR_PTR:.*]] = fir.box_offset %[[DECLARE_2]]#1 base_addr : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
-!HOST: %[[VAR_PTR:.*]] = fir.box_addr %[[LOAD_3]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>) -> !fir.heap<!fir.array<?xi32>>
-!HOST: %[[MAP_INFO_MEMBER:.*]] = omp.map_info var_ptr(%[[VAR_PTR]] : !fir.heap<!fir.array<?xi32>>, !fir.array<?xi32>) var_ptr_ptr(%[[VAR_PTR_PTR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS_2]]) -> !fir.heap<!fir.array<?xi32>> {name = "sp_write(2:5)"}
-!HOST: %[[MAP_INFO_2:.*]] = omp.map_info var_ptr(%[[DECLARE_2]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, !fir.box<!fir.heap<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_INFO_MEMBER]] : !fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> {name = "sp_write(2:5)"}
+!HOST: %[[VAR_PTR:.*]] = fir.box_offset %[[DECLARE_2]]#1 base_addr : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
+!HOST: %[[MAP_INFO_MEMBER:.*]] = omp.map_info var_ptr(%[[VAR_PTR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>, !fir.array<?xi32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS_2]]) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>> {name = ""}
+!HOST: %[[MAP_INFO_2:.*]] = omp.map_info var_ptr(%[[DECLARE_2]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, !fir.box<!fir.heap<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_INFO_MEMBER]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> {name = "sp_write(2:5)"}
subroutine read_write_section()
integer, allocatable :: sp_read(:)
@@ -81,12 +79,10 @@ module assumed_allocatable_array_routines
!HOST: %[[CONSTANT_3:.*]] = arith.constant 0 : index
!HOST: %[[BOX_3:.*]]:3 = fir.box_dims %[[LOAD_3]], %[[CONSTANT_3]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>, index) -> (index, index, index)
!HOST: %[[BOUNDS:.*]] = omp.bounds lower_bound(%[[LB]] : index) upper_bound(%[[UB]] : index) extent(%[[BOX_3]]#1 : index) stride(%[[BOX_2]]#2 : index) start_idx(%[[BOX_1]]#0 : index) {stride_in_bytes = true}
-!HOST: %[[VAR_PTR_PTR:.*]] = fir.box_offset %[[DECLARE]]#1 base_addr : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
-!HOST: %[[VAR_PTR:.*]] = fir.box_addr %[[LOAD_1]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>) -> !fir.heap<!fir.array<?xi32>>
-!HOST: %[[MAP_INFO_MEMBER:.*]] = omp.map_info var_ptr(%[[VAR_PTR]] : !fir.heap<!fir.array<?xi32>>, !fir.array<?xi32>) var_ptr_ptr(%[[VAR_PTR_PTR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS]]) -> !fir.heap<!fir.array<?xi32>> {name = "arr_read_write(2:5)"}
-!HOST: %[[MAP_INFO:.*]] = omp.map_info var_ptr(%[[DECLARE]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, !fir.box<!fir.heap<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_INFO_MEMBER]] : !fir.heap<!fir.array<?xi32>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> {name = "arr_read_write(2:5)"}
-
- subroutine assumed_shape_array(arr_read_write)
+!HOST: %[[VAR_PTR:.*]] = fir.box_offset %[[DECLARE]]#1 base_addr : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
+!HOST: %[[MAP_INFO_MEMBER:.*]] = omp.map_info var_ptr(%[[VAR_PTR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>, !fir.array<?xi32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS]]) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>> {name = ""}
+!HOST: %[[MAP_INFO:.*]] = omp.map_info var_ptr(%[[DECLARE]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, !fir.box<!fir.heap<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_INFO_MEMBER]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> {name = "arr_read_write(2:5)"}
+subroutine assumed_shape_array(arr_read_write)
integer, allocatable, intent(inout) :: arr_read_write(:)
!$omp target map(tofrom:arr_read_write(2:5))
diff --git a/flang/test/Lower/OpenMP/allocatable-map.f90 b/flang/test/Lower/OpenMP/allocatable-map.f90
index 5b8302459ca124..f17c0385e3348e 100644
--- a/flang/test/Lower/OpenMP/allocatable-map.f90
+++ b/flang/test/Lower/OpenMP/allocatable-map.f90
@@ -1,12 +1,10 @@
!RUN: %flang_fc1 -emit-hlfir -fopenmp %s -o - | FileCheck %s --check-prefixes="HLFIRDIALECT"
!HLFIRDIALECT: %[[POINTER:.*]]:2 = hlfir.declare %0 {fortran_attrs = #fir.var_attrs<pointer>, uniq_name = "_QFpointer_routineEpoint"} : (!fir.ref<!fir.box<!fir.ptr<i32>>>) -> (!fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.ref<!fir.box<!fir.ptr<i32>>>)
-!HLFIRDIALECT: %[[LOAD:.*]] = fir.load %[[POINTER]]#1 : !fir.ref<!fir.box<!fir.ptr<i32>>>
!HLFIRDIALECT: %[[BOX_OFF:.*]] = fir.box_offset %[[POINTER]]#1 base_addr : (!fir.ref<!fir.box<!fir.ptr<i32>>>) -> !fir.llvm_ptr<!fir.ref<i32>>
-!HLFIRDIALECT: %[[BOX_ADDR:.*]] = fir.box_addr %[[LOAD]] : (!fir.box<!fir.ptr<i32>>) -> !fir.ptr<i32>
-!HLFIRDIALECT: %[[POINTER_MAP_MEMBER:.*]] = omp.map_info var_ptr(%[[BOX_ADDR]] : !fir.ptr<i32>, i32) var_ptr_ptr(%[[BOX_OFF]] : !fir.llvm_ptr<!fir.ref<i32>>) map_clauses(implicit, tofrom) capture(ByRef) -> !fir.ptr<i32> {name = "point"}
-!HLFIRDIALECT: %[[POINTER_MAP:.*]] = omp.map_info var_ptr(%[[POINTER]]#1 : !fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.box<!fir.ptr<i32>>) map_clauses(implicit, tofrom) capture(ByRef) members(%[[POINTER_MAP_MEMBER]] : !fir.ptr<i32>) -> !fir.ref<!fir.box<!fir.ptr<i32>>> {name = "point"}
-!HLFIRDIALECT: omp.target map_entries({{.*}}, %[[POINTER_MAP]] -> {{.*}} : {{.*}}, !fir.ref<!fir.box<!fir.ptr<i32>>>) {
+!HLFIRDIALECT: %[[POINTER_MAP_MEMBER:.*]] = omp.map_info var_ptr(%[[BOX_OFF]] : !fir.llvm_ptr<!fir.ref<i32>>, i32) map_clauses(implicit, tofrom) capture(ByRef) -> !fir.llvm_ptr<!fir.ref<i32>> {name = ""}
+!HLFIRDIALECT: %[[POINTER_MAP:.*]] = omp.map_info var_ptr(%[[POINTER]]#1 : !fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.box<!fir.ptr<i32>>) map_clauses(implicit, tofrom) capture(ByRef) members(%[[POINTER_MAP_MEMBER]] : !fir.llvm_ptr<!fir.ref<i32>>) -> !fir.ref<!fir.box<!fir.ptr<i32>>> {name = "point"}
+!HLFIRDIALECT: omp.target map_entries({{.*}}, %[[POINTER_MAP_MEMBER]] -> {{.*}}, %[[POINTER_MAP]] -> {{.*}} : {{.*}}, !fir.llvm_ptr<!fir.ref<i32>>, !fir.ref<!fir.box<!fir.ptr<i32>>>) {
subroutine pointer_routine()
integer, pointer :: point
!$omp target map(tofrom:pointer)
diff --git a/flang/test/Lower/OpenMP/array-bounds.f90 b/flang/test/Lower/OpenMP/array-bounds.f90
index 34d32468679b68..07d23233bf7e14 100644
--- a/flang/test/Lower/OpenMP/array-bounds.f90
+++ b/flang/test/Lower/OpenMP/array-bounds.f90
@@ -50,11 +50,10 @@ module assumed_array_routines
!HOST: %[[C0_1:.*]] = arith.constant 0 : index
!HOST: %[[DIMS1:.*]]:3 = fir.box_dims %[[ARG0_DECL]]#1, %[[C0_1]] : (!fir.box<!fir.array<?xi32>>, index) -> (index, index, index)
!HOST: %[[BOUNDS:.*]] = omp.bounds lower_bound(%[[C3]] : index) upper_bound(%[[C4]] : index) extent(%[[DIMS1]]#1 : index) stride(%[[DIMS0]]#2 : index) start_idx(%[[C0]] : index) {stride_in_bytes = true}
-!HOST: %[[VAR_PTR_PTR:.*]] = fir.box_offset %0 base_addr : (!fir.ref<!fir.box<!fir.array<?xi32>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
-!HOST: %[[VAR_PTR:.*]] = fir.box_addr %[[ARG0_DECL]]#1 : (!fir.box<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>>
-!HOST: %[[MAP_INFO_MEMBER:.*]] = omp.map_info var_ptr(%[[VAR_PTR]] : !fir.ref<!fir.array<?xi32>>, !fir.array<?xi32>) var_ptr_ptr(%[[VAR_PTR_PTR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS]]) -> !fir.ref<!fir.array<?xi32>> {name = "arr_read_write(2:5)"}
-!HOST: %[[MAP:.*]] = omp.map_info var_ptr(%[[INTERMEDIATE_ALLOCA]] : !fir.ref<!fir.box<!fir.array<?xi32>>>, !fir.box<!fir.array<?xi32>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_INFO_MEMBER]] : !fir.ref<!fir.array<?xi32>>) -> !fir.ref<!fir.box<!fir.array<?xi32>>> {name = "arr_read_write(2:5)"}
-!HOST: omp.target map_entries(%[[MAP]] -> %{{.*}}, {{.*}} -> {{.*}} : !fir.ref<!fir.box<!fir.array<?xi32>>>, !fir.ref<i32>) {
+!HOST: %[[VAR_PTR:.*]] = fir.box_offset %0 base_addr : (!fir.ref<!fir.box<!fir.array<?xi32>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
+!HOST: %[[MAP_INFO_MEMBER:.*]] = omp.map_info var_ptr(%[[VAR_PTR]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>, !fir.array<?xi32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS]]) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>> {name = ""}
+!HOST: %[[MAP:.*]] = omp.map_info var_ptr(%[[INTERMEDIATE_ALLOCA]] : !fir.ref<!fir.box<!fir.array<?xi32>>>, !fir.box<!fir.array<?xi32>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_INFO_MEMBER]] : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) -> !fir.ref<!fir.array<?xi32>> {name = "arr_read_write(2:5)"}
+!HOST: omp.target map_entries(%[[MAP_INFO_MEMBER]] -> %{{.*}}, %[[MAP]] -> %{{.*}}, {{.*}} -> {{.*}} : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>, !fir.ref<!fir.array<?xi32>>, !fir.ref<i32>) {
subroutine assumed_shape_array(arr_read_write)
integer, intent(inout) :: arr_read_write(:)
diff --git a/flang/test/Lower/OpenMP/target.f90 b/flang/test/Lower/OpenMP/target.f90
index 51571b48645760..11eade38f912f6 100644
--- a/flang/test/Lower/OpenMP/target.f90
+++ b/flang/test/Lower/OpenMP/target.f90
@@ -445,9 +445,9 @@ subroutine omp_target_device_addr
integer, pointer :: a
!CHECK: %[[VAL_0:.*]] = fir.alloca !fir.box<!fir.ptr<i32>> {bindc_name = "a", uniq_name = "_QFomp_target_device_addrEa"}
!CHECK: %[[VAL_0_DECL:.*]]:2 = hlfir.declare %0 {fortran_attrs = #fir.var_attrs<pointer>, uniq_name = "_QFomp_target_device_addrEa"} : (!fir.ref<!fir.box<!fir.ptr<i32>>>) -> (!fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.ref<!fir.box<!fir.ptr<i32>>>)
- !CHECK: %[[MAP_MEMBERS:.*]] = omp.map_info var_ptr({{.*}} : !fir.ptr<i32>, i32) var_ptr_ptr({{.*}} : !fir.llvm_ptr<!fir.ref<i32>>) map_clauses(tofrom) capture(ByRef) -> !fir.ptr<i32> {name = "a"}
- !CHECK: %[[MAP:.*]] = omp.map_info var_ptr({{.*}} : !fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.box<!fir.ptr<i32>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_MEMBERS]] : !fir.ptr<i32>) -> !fir.ref<!fir.box<!fir.ptr<i32>>> {name = "a"}
- !CHECK: omp.target_data map_entries(%[[MAP]] : {{.*}}) use_device_addr(%[[VAL_0_DECL]]#1 : !fir.ref<!fir.box<!fir.ptr<i32>>>) {
+ !CHECK: %[[MAP_MEMBERS:.*]] = omp.map_info var_ptr({{.*}} : !fir.llvm_ptr<!fir.ref<i32>>, i32) map_clauses(tofrom) capture(ByRef) -> !fir.llvm_ptr<!fir.ref<i32>> {name = ""}
+ !CHECK: %[[MAP:.*]] = omp.map_info var_ptr({{.*}} : !fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.box<!fir.ptr<i32>>) map_clauses(tofrom) capture(ByRef) members(%[[MAP_MEMBERS]] : !fir.llvm_ptr<!fir.ref<i32>>) -> !fir.ref<!fir.box<!fir.ptr<i32>>> {name = "a"}
+ !CHECK: omp.target_data map_entries(%[[MAP_MEMBERS]], %[[MAP]] : {{.*}}) use_device_addr(%[[VAL_0_DECL]]#1 : !fir.ref<!fir.box<!fir.ptr<i32>>>) {
!$omp target data map(tofrom: a) use_device_addr(a)
!CHECK: ^bb0(%[[VAL_1:.*]]: !fir.ref<!fir.box<!fir.ptr<i32>>>):
!CHECK: %[[VAL_1_DECL:.*]]:2 = hlfir.declare %[[VAL_1]] {fortran_attrs = #fir.var_attrs<pointer>, uniq_name = "_QFomp_target_device_addrEa"} : (!fir.ref<!fir.box<!fir.ptr<i32>>>) -> (!fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.ref<!fir.box<!fir.ptr<i32>>>)
diff --git a/mlir/test/Target/LLVMIR/omptarget-fortran-allocatable-types-host.mlir b/mlir/test/Target/LLVMIR/omptarget-fortran-allocatable-types-host.mlir
index dc3b230a7a55ee..71c49102f775b4 100644
--- a/mlir/test/Target/LLVMIR/omptarget-fortran-allocatable-types-host.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-fortran-allocatable-types-host.mlir
@@ -118,8 +118,8 @@ module attributes {omp.is_target_device = false} {
%95 = llvm.getelementptr %14[0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
%96 = omp.map_info var_ptr(%95 : !llvm.ptr, f32) map_clauses(tofrom) capture(ByRef) -> !llvm.ptr {name = "scalar"}
%97 = omp.map_info var_ptr(%14 : !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>) map_clauses(tofrom) capture(ByRef) members(%96 : !llvm.ptr) -> !llvm.ptr {name = "scalar"}
- omp.target map_entries(%81 -> %arg0, %94 -> %arg1, %97 -> %arg2 : !llvm.ptr, !llvm.ptr, !llvm.ptr) {
- ^bb0(%arg0: !llvm.ptr, %arg1: !llvm.ptr, %arg2: !llvm.ptr):
+ omp.target map_entries(%80 -> %arg0, %81 -> %arg1, %93 -> %arg2, %94 -> %arg3, %96 -> %arg4, %97 -> %arg5 : !llvm.ptr, !llvm.ptr, !llvm.ptr, !llvm.ptr, !llvm.ptr, !llvm.ptr) {
+ ^bb0(%arg0: !llvm.ptr, %arg1: !llvm.ptr, %arg2: !llvm.ptr, %arg3: !llvm.ptr, %arg4: !llvm.ptr, %arg5: !llvm.ptr):
%98 = llvm.mlir.constant(1 : i32) : i32
%99 = llvm.alloca %98 x !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)> {alignment = 8 : i64} : (i32) -> !llvm.ptr
%100 = llvm.alloca %98 x !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)> {alignment = 8 : i64} : (i32) -> !llvm.ptr
@@ -129,7 +129,7 @@ module attributes {omp.is_target_device = false} {
%104 = llvm.mlir.constant(3 : index) : i64
%105 = llvm.mlir.constant(1 : index) : i64
%106 = llvm.mlir.constant(1.000000e+00 : f32) : f32
- %107 = llvm.load %arg0 : !llvm.ptr -> !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %107 = llvm.load %arg1 : !llvm.ptr -> !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
llvm.store %107, %101 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>, !llvm.ptr
%108 = llvm.getelementptr %101[0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
%109 = llvm.load %108 : !llvm.ptr -> !llvm.ptr
@@ -146,7 +146,7 @@ module attributes {omp.is_target_device = false} {
%120 = llvm.add %119, %116 : i64
%121 = llvm.getelementptr %109[%120] : (!llvm.ptr, i64) -> !llvm.ptr, f32
llvm.store %106, %121 : f32, !llvm.ptr
- %122 = llvm.load %arg1 : !llvm.ptr -> !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
+ %122 = llvm.load %arg3 : !llvm.ptr -> !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
llvm.store %122, %100 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>, !llvm.ptr
%123 = llvm.getelementptr %100[0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<1 x array<3 x i64>>)>
%124 = llvm.load %123 : !llvm.ptr -> !llvm.ptr
@@ -182,7 +182,7 @@ module attributes {omp.is_target_device = false} {
%152 = llvm.insertvalue %136, %151[0] : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>
llvm.store %152, %99 : !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8)>, !llvm.ptr
%153 = llvm.mlir.addressof @_QQclX6606f061cfd55e600f1165aa97a647a5 : !llvm.ptr
- %154 = llvm.call @_FortranAAssign(%arg2, %99, %153, %102) : (!llvm.ptr, !llvm.ptr, !llvm.ptr, i32) -> !llvm.struct<()>
+ %154 = llvm.call @_FortranAAssign(%arg5, %99, %153, %102) : (!llvm.ptr, !llvm.ptr, !llvm.ptr, i32) -> !llvm.struct<()>
omp.terminator
}
llvm.return
>From fb46edc7aaa1f0f2e615ccc24c90f4d0c8cdc972 Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Fri, 5 Jan 2024 09:23:21 -0600
Subject: [PATCH 18/19] [Documentation][NFC] Add document describing descriptor
mapping in a little bit of detail, and remove some unneccesary comments from
new pass
---
flang/docs/OpenMP-descriptor-management.md | 95 +++++++++++++++++++
.../Transforms/OMPDescriptorMapInfoGen.cpp | 4 -
2 files changed, 95 insertions(+), 4 deletions(-)
create mode 100644 flang/docs/OpenMP-descriptor-management.md
diff --git a/flang/docs/OpenMP-descriptor-management.md b/flang/docs/OpenMP-descriptor-management.md
new file mode 100644
index 00000000000000..ea364ec4dc6d0e
--- /dev/null
+++ b/flang/docs/OpenMP-descriptor-management.md
@@ -0,0 +1,95 @@
+<!--===- docs/OpenMP-descriptor-management.md
+
+ Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+ See https://llvm.org/LICENSE.txt for license information.
+ SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+-->
+
+# OpenMP dialect: Fortran descriptor type mapping for offload
+
+The descriptor mapping for OpenMP currently works differently to the planned direction for OpenACC, however,
+it is possible and would likely be ideal to align the method with OpenACC in the future. However, at least
+currently the OpenMP specification is less descriptive and has less stringent rules around descriptor based
+types so does not require as complex a set of descriptor management rules (although, in certain cases
+for the interim adopting OpenACC's rules where it makes sense could be useful).
+
+The initial method for mapping Fortran types tied to descriptors for OpenMP offloading is to treat these types
+as a special case of OpenMP record type (C/C++ structure/class, Fortran derived type etc.) mapping as far as the
+runtime is concerned. Where the box (descriptor information) is the holding container and the underlying
+data pointer is contained within the container, and we must generate explicit maps for both the pointer member and
+the container. As an example, a small C++ program that is equivalent to the concept described:
+
+```C++
+struct mock_descriptor {
+ long int x;
+ std::byte x1, x2, x3, x4;
+ void *pointer;
+ long int lx[1][3];
+};
+
+int main() {
+mock_descriptor data;
+#pragma omp target map(tofrom: data, data.pointer[:upper_bound])
+{
+ do something...
+}
+
+ return 0;
+}
+```
+
+In the above, we have to map both the containing structure, with its non-pointer members and the
+data pointed to by the pointer contained within the structure to appropriately access the data. This
+is effectively what is done with descriptor types for the time being. Other pointers that are part
+of the descriptor container such as the addendum should also be treated as the data pointer is
+treated.
+
+Currently, Flang will lower these descriptor types in the OpenMP lowering (lower/OpenMP.cpp) similarly
+to all other map types, generating an omp.MapInfoOp containing relevant information required for lowering
+the OpenMP dialect to LLVM-IR during the final stages of the MLIR lowering. However, after
+the lowering to FIR/HLFIR has been performed an OpenMP dialect specific pass for Fortran,
+OMPDescriptorMapInfoGenPass (Optimizer/OMPDescriptorMapInfoGen.cpp) will expand the
+omp.MapInfoOp's containing descriptors (which currently will be a BoxType or BoxAddrOp) into multiple
+mappings, with one extra per pointer member in the descriptor that is supported on top of the original
+descriptor map operation. These pointers members are linked to the parent descriptor by adding them to
+the member field of the original descriptor map operation, they are then inserted into the relevant map
+owning operation's (omp.TargetOp, omp.DataOp etc.) map operand list and in cases where the owning operation
+is IsolatedFromAbove, it also inserts them as BlockArgs to canonicalize the mappings and simplify lowering.
+
+An example transformation by the OMPDescriptorMapInfoGenPass:
+
+```
+
+...
+%12 = omp.map_info var_ptr(%1#1 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.box<!fir.ptr<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) bounds(%11) -> !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>> {name = "arg_alloc"}
+...
+omp.target map_entries(%12 -> %arg1, %13 -> %arg2 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.ref<i32>) {
+ ^bb0(%arg1: !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, %arg2: !fir.ref<i32>):
+...
+
+====>
+
+...
+%12 = fir.box_offset %1#1 base_addr : (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>
+%13 = omp.map_info var_ptr(%12 : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>, !fir.array<?xi32>) map_clauses(tofrom) capture(ByRef) bounds(%11) -> !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>> {name = ""}
+%14 = omp.map_info var_ptr(%1#1 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.box<!fir.ptr<!fir.array<?xi32>>>) map_clauses(tofrom) capture(ByRef) members(%13 : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>) -> !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>> {name = "arg_alloc"}
+...
+omp.target map_entries(%13 -> %arg1, %14 -> %arg2, %15 -> %arg3 : !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>, !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.ref<i32>) {
+ ^bb0(%arg1: !fir.llvm_ptr<!fir.ref<!fir.array<?xi32>>>, %arg2: !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, %arg3: !fir.ref<i32>):
+...
+
+```
+
+In later stages of the compilation flow when the OpenMP dialect is being lowered to LLVM-IR these descriptor
+mappings are treated as if they were structure mappings with explicit member maps on the same directive as
+their parent was mapped.
+
+This method is generic in the sense that the OpenMP diaelct doesn't need to understand that it is mapping a
+Fortran type containing a descriptor, it just thinks it's a record type from either Fortran or C++. However,
+it is a little rigid in how the descriptor mappings are handled as there is no specialisation or possibility
+to specialise the mappings for possible edge cases without poluting the dialect or lowering with further
+knowledge of Fortran and the FIR dialect. In the case that this kind of specialisation is required or
+desired then the methodology described by OpenACC which utilises runtime functions to handle specialised mappings
+for dialects may be a more desirable approach to move towards. For the moment this method appears sufficient as
+far as the OpenMP specification and current testing can show.
diff --git a/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp b/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp
index 17e0ce05143e03..203fdc8ce21ee8 100644
--- a/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp
+++ b/flang/lib/Optimizer/Transforms/OMPDescriptorMapInfoGen.cpp
@@ -80,7 +80,6 @@ class OMPDescriptorMapInfoGenPass
op.getMembersMutable().append(descriptorBaseAddrMembers);
op.getBoundsMutable().assign(llvm::SmallVector<mlir::Value>{});
- // could use a template to generalise to other TargetOps
if (auto mapClauseOwner =
llvm::dyn_cast<mlir::omp::MapClauseOwningOpInterface>(target)) {
llvm::SmallVector<mlir::Value> newMapOps;
@@ -117,9 +116,6 @@ class OMPDescriptorMapInfoGenPass
fir::KindMapping kindMap = fir::getKindMapping(getOperation());
fir::FirOpBuilder builder{getOperation(), std::move(kindMap)};
- // 2) look into adding some documentation on this.
- // 3) fix tests.
-
getOperation()->walk([&](mlir::omp::MapInfoOp op) {
if (fir::isTypeWithDescriptor(op.getVarType()) ||
mlir::isa_and_present<fir::BoxAddrOp>(
>From fe4ce300ebc9532f8abea1a604ef476a37ace693 Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Fri, 5 Jan 2024 10:10:33 -0600
Subject: [PATCH 19/19] remove debug artifact I added
---
mlir/lib/Target/LLVMIR/ModuleTranslation.cpp | 2 --
1 file changed, 2 deletions(-)
diff --git a/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp b/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
index d6c6c7f86f85c4..1722d74c08b628 100644
--- a/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
@@ -1584,8 +1584,6 @@ mlir::translateModuleToLLVMIR(Operation *module, llvm::LLVMContext &llvmContext,
if (failed(translator.convertFunctions()))
return nullptr;
- // translator.llvmModule->dump();
-
if (llvm::verifyModule(*translator.llvmModule, &llvm::errs()))
return nullptr;
More information about the flang-commits
mailing list