[flang-commits] [flang] [Flang][Lower][OpenMP] Add initial lowering of pointers/allocatables/target in map clauses to map_info and entries (PR #68600)
via flang-commits
flang-commits at lists.llvm.org
Mon Oct 9 08:42:04 PDT 2023
https://github.com/agozillon created https://github.com/llvm/llvm-project/pull/68600
This patch seeks to add an initial lowering for pointers, allocatables and target variables explicitly captured by map in Flang OpenMP.
Currently the intention is to treat these types as a special case of OpenMP structure mapping as far as the runtime is concerned, where the box (descriptor information) is the holding container and the underlying pointer is contained within the container. The descriptor pointed to by the generated bounds provides all the data required to offload the pointer contained within the descriptor.
This comes from the concept that the eventual lowered LLVM IR Type for these types is effectively a structure containing descriptor information and a pointer to the data. And the target device kernel essentially treats these as such.
A future optimization might be to strip the descriptor information to turn these into naked pointers, but I feel the descriptor information is rather important and perhaps the OpenMP specification has something to say about maintaining Fortran descriptors!
This patch generates two map_info for each Fortran pointer, target or allocatable, implicitly mimicking the following relationship in C++:
template<typename T>
struct descriptor {
... other descriptor data
T* data;
};
#pragma omp target map (tofrom: descriptor, descriptor->data)
Or in Fortran:
type :: descriptor
... other descriptor data
integer :: data(:) ! not completely accurate but showcases the idea
end type scalar
!$omp target map(from:scalar, scalar%data)
Where the first map generated is the descriptor and the second argument is the pointer to the data. This results in the following operations generated per mapped allocatable/target/pointer:
1) omp.bounds pointing to the descriptor information relevant to the
bounds of the described data
2) one map which contains no bounds information and it's varptr (main
map argument) points to the descriptor (box), this is the first map
in the above examples indicating the descriptor container.
3) a second map which contains the bounds operation from 1, a varptr
pointing to the data/address of the to be mapped
pointer/target/allocatable variable contained within the descriptor
(we generate a box_addr) and then a subsequent varptptr (which
points to an owning object) which points to the descriptor, that
was mapped in the first map, indicating a link between these two
maps for later lowering to take advantage of.
NOTE: Currently there's a block on pointers/allocas/target contained in derived types
as it appears the bounds information isn't fully generated for these just yet.
>From 873c181efe0d813c79c4f5d96b68b14469a95a7d Mon Sep 17 00:00:00 2001
From: Andrew Gozillon <Andrew.Gozillon at amd.com>
Date: Mon, 9 Oct 2023 10:33:10 -0500
Subject: [PATCH] [Flang][Lower][OpenMP] Add initial lowering of
pointers/allocatables/target in map clauses to map_info and entries
This patch seeks to add an initial lowering for pointers, allocatables
and target variables explicitly captured by map in Flang OpenMP.
Currently the intention is to treat these types as a special case of
OpenMP structure mapping as far as the runtime is concerned,
where the box (descriptor information) is the holding container
and the underlying pointer is contained within the container.
The descriptor pointed to by the generated bounds provides
all the data required to offload the pointer contained within
the descriptor.
This comes from the concept that the eventual lowered
LLVM IR Type for these types is effectively a structure
containing descriptor information and a pointer to
the data. And the target device kernel essentially
treats these as such.
A future optimisation might be to strip the descriptor
information to turn these into naked pointers, but I
feel the descriptor information is rather important
and perhaps the OpenMP specification has something
to say about maintaining Fortran descriptors!
This patch generates two map_info for each Fortran
pointer, target or allocatable, implicitly mimicing the
following relationship in C++:
template<typename T>
struct descriptor {
... other descriptor data
T* data;
};
#pragma omp target map (tofrom: descriptor, descriptor->data)
Or in Fortran:
type :: descriptor
... other descriptor data
integer :: data(:) ! not completely accurate but showcases the idea
end type scalar
!$omp target map(from:scalar, scalar%data)
Where the first map generated is the descriptor and the second argument
is the pointer to the data. This results in the following operations generated
per mapped allocatable/target/pointer:
1) omp.bounds pointing to the descriptor information relevant to the
bounds of the described data
2) one map which contains no bounds information and it's varptr (main
map argument) points to the descriptor (box), this is the first map
in the above examples indicating the descriptor container.
3) a second map which contains the bounds operation from 1, a varptr
pointing to the data/address of the to be mapped
pointer/target/allocatable variable contained within the descriptor
(we generate a box_addr) and then a subsequent varptptr (which
points to an owning object) which points to the descriptor, that
was mapped in the first map, indicating a link between these two
maps for later lowering to take adavantage of.
---
flang/lib/Lower/OpenMP.cpp | 87 +++++++++++++++++--
flang/test/Lower/OpenMP/FIR/target.f90 | 5 +-
.../OpenMP/allocatable-pointer-target-map.f90 | 74 ++++++++++++++++
3 files changed, 157 insertions(+), 9 deletions(-)
create mode 100644 flang/test/Lower/OpenMP/allocatable-pointer-target-map.f90
diff --git a/flang/lib/Lower/OpenMP.cpp b/flang/lib/Lower/OpenMP.cpp
index 5f5e968eaaa6414..8002423e0a1961c 100644
--- a/flang/lib/Lower/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP.cpp
@@ -48,6 +48,9 @@ getOmpObjectSymbol(const Fortran::parser::OmpObject &ompObject) {
Fortran::parser::Unwrap<Fortran::parser::ArrayElement>(
designator)) {
sym = GetFirstName(arrayEle->base).symbol;
+ } else if (auto *structComp = Fortran::parser::Unwrap<
+ Fortran::parser::StructureComponent>(designator)) {
+ sym = structComp->component.symbol;
} else if (const Fortran::parser::Name *name =
Fortran::semantics::getDesignatorNameIfDataRef(
designator)) {
@@ -1663,11 +1666,10 @@ 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, bool implicit,
- mlir::Type retTy) {
- mlir::Value varPtrPtr;
+ mlir::Value baseAddr, mlir::Value varPtrPtr,
+ const std::string &name, mlir::SmallVector<mlir::Value> bounds,
+ uint64_t mapType, mlir::omp::VariableCaptureKind mapCaptureType,
+ bool implicit, mlir::Type retTy) {
if (auto boxTy = baseAddr.getType().dyn_cast<fir::BaseBoxType>()) {
baseAddr = builder.create<fir::BoxAddrOp>(loc, baseAddr);
retTy = baseAddr.getType();
@@ -1675,7 +1677,7 @@ createMapInfoOp(fir::FirOpBuilder &builder, mlir::Location loc,
mlir::omp::MapInfoOp op =
builder.create<mlir::omp::MapInfoOp>(loc, retTy, baseAddr);
- op.setNameAttr(builder.getStringAttr(name.str()));
+ op.setNameAttr(builder.getStringAttr(name));
op.setImplicit(implicit);
op.setMapType(mapType);
op.setMapCaptureType(mapCaptureType);
@@ -1752,11 +1754,82 @@ bool ClauseProcessor::processMap(
semanticsContext, stmtCtx, ompObject,
clauseLocation, asFortran, bounds);
+ auto checkIfStructComponent = [](const Fortran::parser::OmpObject
+ &ompObject) {
+ bool isComponent = false;
+ std::visit(
+ Fortran::common::visitors{
+ [&](const Fortran::parser::Designator &designator) {
+ if (auto *structComp = Fortran::parser::Unwrap<
+ Fortran::parser::StructureComponent>(
+ designator)) {
+ if (std::holds_alternative<Fortran::parser::Name>(
+ structComp->base.u))
+ isComponent = true;
+ }
+ },
+ [&](const Fortran::parser::Name &name) {}},
+ ompObject.u);
+
+ return isComponent;
+ };
+
+ // TODO: Currently, it appears there's missing symbol information
+ // and bounds information for allocatables and pointers inside
+ // of derived types. The latter needs some additional support
+ // added to the bounds generation whereas the former appears
+ // that it could be a problem when referring to pointer members
+ // via an OpenMP map clause, for the moment we do not handle
+ // these cases and must emit an error.
+ if (checkIfStructComponent(ompObject) &&
+ Fortran::semantics::IsAllocatableOrPointer(
+ *getOmpObjectSymbol(ompObject)))
+ TODO(currentLocation,
+ "pointer members of derived types are currently unmapped");
+
+ if (Fortran::semantics::IsAllocatableOrPointer(
+ *getOmpObjectSymbol(ompObject))) {
+ // We mimic what will eventually be a structure containing a
+ // pointer mapping for allocatables/pointers/target e.g.:
+ //
+ // !$omp target map(from:in, in%map_ptr)
+ //
+ // ===>
+ //
+ // map_entry varptr(in) ....
+ // map_entry varptr(map_ptr) varptrptr(in) ...
+ //
+ // This is to attempt to keep the lowering of these consistent
+ // with structures containing pointers that are mapped like the
+ // example above, where we break it into the descriptor being the
+ // main "structure" being mapped and the contained pointer the
+ // specific member being referenced. This is of course implicit,
+ // the user just maps the pointer, target or allocatable.
+ mlir::Value descriptor =
+ converter.getSymbolAddress(*getOmpObjectSymbol(ompObject));
+ mapOperands.push_back(createMapInfoOp(
+ firOpBuilder, clauseLocation, descriptor, nullptr,
+ asFortran.str(), mlir::SmallVector<mlir::Value>{},
+ static_cast<std::underlying_type_t<
+ llvm::omp::OpenMPOffloadMappingFlags>>(mapTypeBits),
+ mlir::omp::VariableCaptureKind::ByRef, false,
+ descriptor.getType()));
+ mapOperands.push_back(createMapInfoOp(
+ firOpBuilder, clauseLocation, baseAddr, descriptor,
+ asFortran.str(), bounds,
+ static_cast<std::underlying_type_t<
+ llvm::omp::OpenMPOffloadMappingFlags>>(mapTypeBits),
+ mlir::omp::VariableCaptureKind::ByRef, false,
+ baseAddr.getType()));
+ continue;
+ }
+
// Explicit map captures are captured ByRef by default,
// optimisation passes may alter this to ByCopy or other capture
// types to optimise
mapOperands.push_back(createMapInfoOp(
- firOpBuilder, clauseLocation, baseAddr, asFortran, bounds,
+ firOpBuilder, clauseLocation, baseAddr, nullptr, asFortran.str(),
+ bounds,
static_cast<
std::underlying_type_t<llvm::omp::OpenMPOffloadMappingFlags>>(
mapTypeBits),
diff --git a/flang/test/Lower/OpenMP/FIR/target.f90 b/flang/test/Lower/OpenMP/FIR/target.f90
index 9b1fb5c15ac1d2d..6267bb4221a095c 100644
--- a/flang/test/Lower/OpenMP/FIR/target.f90
+++ b/flang/test/Lower/OpenMP/FIR/target.f90
@@ -248,8 +248,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:.*]] = omp.map_info var_ptr({{.*}}) map_clauses(tofrom) capture(ByRef) -> {{.*}} {name = "a"}
- !CHECK: omp.target_data map_entries(%[[MAP]] : {{.*}}) use_device_addr(%[[VAL_0]] : !fir.ref<!fir.box<!fir.ptr<i32>>>) {
+ !CHECK: %[[MAP_DESC:.*]] = omp.map_info var_ptr({{.*}}) map_clauses(tofrom) capture(ByRef) -> {{.*}} {name = "a"}
+ !CHECK: %[[MAP_PTR:.*]] = omp.map_info var_ptr({{.*}}) var_ptr_ptr({{.*}}) map_clauses(tofrom) capture(ByRef) -> {{.*}} {name = "a"}
+ !CHECK: omp.target_data map_entries(%[[MAP_DESC]], %[[MAP_PTR]] : {{.*}}) 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-pointer-target-map.f90 b/flang/test/Lower/OpenMP/allocatable-pointer-target-map.f90
new file mode 100644
index 000000000000000..7cbe0e5db4d5f66
--- /dev/null
+++ b/flang/test/Lower/OpenMP/allocatable-pointer-target-map.f90
@@ -0,0 +1,74 @@
+
+!RUN: %flang_fc1 -emit-hlfir -fopenmp %s -o - | FileCheck %s
+
+subroutine map_pointer()
+ integer, pointer :: map_ptr(:)
+ allocate(map_ptr(10))
+ !CHECK: %[[ALLOCA:.*]] = fir.alloca !fir.box<!fir.ptr<!fir.array<?xi32>>> {bindc_name = "map_ptr", uniq_name = "_QFmap_pointerEmap_ptr"}
+ !CHECK: %[[DESC:.*]]:2 = hlfir.declare %[[ALLOCA]] {fortran_attrs = #fir.var_attrs<pointer>, uniq_name = "_QFmap_pointerEmap_ptr"} : (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>) -> (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>)
+ !CHECK: %[[LOAD_FROM_DESC:.*]] = fir.load %[[DESC]]#1 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>
+ !CHECK: %[[MAP_BOUNDS:.*]] = omp.bounds lower_bound({{.*}}) upper_bound({{.*}}) stride({{.*}}) start_idx({{.*}}) {stride_in_bytes = true}
+ !CHECK: %[[MAP_DESC:.*]] = omp.map_info var_ptr(%[[DESC]]#1 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>) map_clauses(tofrom) capture(ByRef) -> !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>> {name = "map_ptr"}
+ !CHECK: %[[PTR_ADDR:.*]] = fir.box_addr %[[LOAD_FROM_DESC]] : (!fir.box<!fir.ptr<!fir.array<?xi32>>>) -> !fir.ptr<!fir.array<?xi32>>
+ !CHECK: %[[MAP_PTR:.*]] = omp.map_info var_ptr(%[[PTR_ADDR]] : !fir.ptr<!fir.array<?xi32>>) var_ptr_ptr(%[[DESC]]#1 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>) map_clauses(tofrom) capture(ByRef) bounds(%[[MAP_BOUNDS]]) -> !fir.ptr<!fir.array<?xi32>> {name = "map_ptr"}
+ !CHECK: omp.target map_entries(%[[MAP_DESC]], %[[MAP_PTR]] : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.ptr<!fir.array<?xi32>>) {
+ !$omp target map(tofrom: map_ptr)
+ !$omp end target
+end subroutine map_pointer
+
+subroutine map_alloca()
+ integer, allocatable :: map_al(:)
+ allocate(map_al(10))
+ !CHECK: %[[ALLOCA:.*]] = fir.alloca !fir.box<!fir.heap<!fir.array<?xi32>>> {bindc_name = "map_al", uniq_name = "_QFmap_allocaEmap_al"}
+ !CHECK: %[[DESC:.*]]:2 = hlfir.declare %[[ALLOCA]] {fortran_attrs = #fir.var_attrs<allocatable>, uniq_name = "_QFmap_allocaEmap_al"} : (!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>>>>)
+ !CHECK: %[[LOAD_FROM_DESC:.*]] = fir.load %[[DESC]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>
+ !CHECK: %[[MAP_BOUNDS:.*]] = omp.bounds lower_bound({{.*}}) upper_bound({{.*}}) stride({{.*}}) start_idx({{.*}}) {stride_in_bytes = true}
+ !CHECK: %[[MAP_DESC:.*]] = omp.map_info var_ptr(%[[DESC]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) map_clauses(tofrom) capture(ByRef) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> {name = "map_al"}
+ !CHECK: %[[PTR_ADDR:.*]] = fir.box_addr %[[LOAD_FROM_DESC]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>) -> !fir.heap<!fir.array<?xi32>>
+ !CHECK: %[[MAP_PTR:.*]] = omp.map_info var_ptr(%[[PTR_ADDR]] : !fir.heap<!fir.array<?xi32>>) var_ptr_ptr(%[[DESC]]#1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) map_clauses(tofrom) capture(ByRef) bounds(%[[MAP_BOUNDS]]) -> !fir.heap<!fir.array<?xi32>> {name = "map_al"}
+ !CHECK: omp.target map_entries(%[[MAP_DESC]], %[[MAP_PTR]] : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, !fir.heap<!fir.array<?xi32>>) {
+ !$omp target map(tofrom: map_al)
+ !$omp end target
+end subroutine map_alloca
+
+subroutine map_pointer_target()
+ integer, pointer :: a(:)
+ integer, target :: b(10)
+ a => b
+ !CHECK: %[[ALLOCA:.*]] = fir.alloca !fir.box<!fir.ptr<!fir.array<?xi32>>> {bindc_name = "a", uniq_name = "_QFmap_pointer_targetEa"}
+ !CHECK: %[[DESC:.*]]:2 = hlfir.declare %[[ALLOCA]] {fortran_attrs = #fir.var_attrs<pointer>, uniq_name = "_QFmap_pointer_targetEa"} : (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>) -> (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>)
+ !CHECK: %[[LOAD_FROM_DESC:.*]] = fir.load %[[DESC]]#1 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>
+ !CHECK: %[[MAP_BOUNDS:.*]] = omp.bounds lower_bound({{.*}}) upper_bound({{.*}}) stride({{.*}}) start_idx({{.*}}) {stride_in_bytes = true}
+ !CHECK: %[[MAP_DESC:.*]] = omp.map_info var_ptr(%[[DESC]]#1 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>) map_clauses(tofrom) capture(ByRef) -> !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>> {name = "a"}
+ !CHECK: %[[PTR_ADDR:.*]] = fir.box_addr %[[LOAD_FROM_DESC]] : (!fir.box<!fir.ptr<!fir.array<?xi32>>>) -> !fir.ptr<!fir.array<?xi32>>
+ !CHECK: %[[MAP_PTR:.*]] = omp.map_info var_ptr(%[[PTR_ADDR]] : !fir.ptr<!fir.array<?xi32>>) var_ptr_ptr(%[[DESC]]#1 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>) map_clauses(tofrom) capture(ByRef) bounds(%[[MAP_BOUNDS]]) -> !fir.ptr<!fir.array<?xi32>> {name = "a"}
+ !CHECK: omp.target map_entries(%[[MAP_DESC]], %[[MAP_PTR]] : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.ptr<!fir.array<?xi32>>) {
+ !$omp target map(tofrom: a)
+ !$omp end target
+end subroutine map_pointer_target
+
+subroutine map_pointer_target_section()
+ integer,target :: A(30)
+ integer,pointer :: p(:)
+ !CHECK: %[[ALLOCA:.*]] = fir.alloca !fir.array<30xi32> {bindc_name = "a", fir.target, uniq_name = "_QFmap_pointer_target_sectionEa"}
+ !CHECK: %[[SHAPE:.*]] = fir.shape %c30 : (index) -> !fir.shape<1>
+ !CHECK: %[[DESC_1:.*]]:2 = hlfir.declare %[[ALLOCA]](%[[SHAPE]]) {fortran_attrs = #fir.var_attrs<target>, uniq_name = "_QFmap_pointer_target_sectionEa"} : (!fir.ref<!fir.array<30xi32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<30xi32>>, !fir.ref<!fir.array<30xi32>>)
+ !CHECK: %[[ALLOCA_2:.*]] = fir.alloca !fir.box<!fir.ptr<!fir.array<?xi32>>> {bindc_name = "p", uniq_name = "_QFmap_pointer_target_sectionEp"}
+ !CHECK: %[[DESC_2:.*]]:2 = hlfir.declare %[[ALLOCA_2]] {fortran_attrs = #fir.var_attrs<pointer>, uniq_name = "_QFmap_pointer_target_sectionEp"} : (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>) -> (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>)
+ !CHECK: %[[MAP_1_BOUNDS:.*]] = omp.bounds lower_bound({{.*}}) upper_bound({{.*}}) stride({{.*}}) start_idx({{.*}})
+ !CHECK: %[[MAP_1:.*]] = omp.map_info var_ptr(%[[DESC_1]]#1 : !fir.ref<!fir.array<30xi32>>) map_clauses(tofrom) capture(ByRef) bounds(%[[MAP_1_BOUNDS]]) -> !fir.ref<!fir.array<30xi32>> {name = "a(1:4)"}
+ !CHECK: omp.target_data map_entries(%[[MAP_1]] : !fir.ref<!fir.array<30xi32>>) {
+ !$omp target data map( A(1:4) )
+ p=>A
+ !CHECK: %[[LOAD:.*]] = fir.load %[[DESC_2]]#1 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>
+ !CHECK: %[[MAP_3_BOUNDS:.*]] = omp.bounds lower_bound({{.*}}) upper_bound({{.*}}) stride({{.*}}) start_idx({{.*}}) {stride_in_bytes = true}
+ !CHECK: %[[MAP_2:.*]] = omp.map_info var_ptr(%[[DESC_2]]#1 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>) map_clauses(tofrom) capture(ByRef) -> !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>> {name = "p(8:27)"}
+ !CHECK: %[[MAP_ADDR_OF:.*]] = fir.box_addr %[[LOAD]] : (!fir.box<!fir.ptr<!fir.array<?xi32>>>) -> !fir.ptr<!fir.array<?xi32>>
+ !CHECK: %[[MAP_3:.*]] = omp.map_info var_ptr(%[[MAP_ADDR_OF]] : !fir.ptr<!fir.array<?xi32>>) var_ptr_ptr(%[[DESC_2]]#1 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>) map_clauses(tofrom) capture(ByRef) bounds(%[[MAP_3_BOUNDS]]) -> !fir.ptr<!fir.array<?xi32>> {name = "p(8:27)"}
+ !CHECK: omp.target map_entries(%[[MAP_2]], %[[MAP_3]] : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, !fir.ptr<!fir.array<?xi32>>) {
+ !$omp target map( p(8:27) )
+ A(3) = 0
+ p(9) = 0
+ !$omp end target
+ !$omp end target data
+end subroutine map_pointer_target_section
More information about the flang-commits
mailing list