From: Julian Brown <julian@codesourcery.com>
To: Tobias Burnus <tobias@codesourcery.com>
Cc: <gcc-patches@gcc.gnu.org>, <fortran@gcc.gnu.org>,
Jakub Jelinek <jakub@redhat.com>
Subject: Re: [PATCH v5 3/4] OpenMP: Pointers and member mappings
Date: Thu, 15 Dec 2022 14:54:58 +0000 [thread overview]
Message-ID: <20221215145458.458074ca@squid.athome> (raw)
In-Reply-To: <aa0da1c7-4bc3-6fc3-ee00-1c2cc04fdcde@codesourcery.com>
[-- Attachment #1: Type: text/plain, Size: 2519 bytes --]
On Wed, 7 Dec 2022 17:31:20 +0100
Tobias Burnus <tobias@codesourcery.com> wrote:
> Hi Julian,
>
> I think this patch is OK; however, at least for gimplify.cc Jakub
> needs to have a second look.
Thanks for the review! Here's a new version that hopefully addresses
your comments. (The gimplify bits change a bit more in this version!)
> As remarked for the 2/4 patch, I believe mapping 'map(tofrom:
> var%f(2:3))' should work without explicitly mapping 'map(tofrom:
> var%f)' (→ [TR11 157:21-26] (approx. [5.2 154:22-27], [5.1
> 352:17-22], [5.0 320:22-27]). →
> https://gcc.gnu.org/pipermail/gcc-patches/2022-December/608100.html
> (+ previously in the thread).
>
> Testing the patch, that seems to work fine (i.e. contrary to C/C++,
> cf. 2/4), which matches the dump and, if I understood correctly, also
> your (Julian's) expectation. Thus, no need to modify the code part.
That change is undone.
> Regarding the testcases:
> * I would prefer if you don't modify the existing
> libgomp.fortran/struct-elem-map-1.f90 testcase; However, you could
> add your version as another variant ('subroutine nine()',
> 'four_var()' or what's the next free name, possibly with a comment
> telling that it is 'four()' but with an added explicit basepointer
> mapping.).
I've added new functions "nine" through "twelve" with the added base
pointer.
> * As the new version should map *less*, I wonder whether some
> -fdump-tree-{original,gimple,omplower} scan-dump-tree checks would be
> useful besides testing whether it works at run time. (Your decision
> regarding which tree, which testcases and whether at all.)
A couple added...
> * Likewise, maybe a 'target enter/exit data' check? However, you
> might very well run into my 'omp target data exit' issue, cf.
> https://gcc.gnu.org/pipermail/gcc-patches/2022-November/604887.html
> (needs to be revised based on Jakub's comments; I think those were on
> IRC only – the problem is that not only 'alloc' is affected but also
> 'from' etc.)
Hmm -- this wasn't quite as straightforward as it probably should have
been! This version of the patch fixes a couple of bugs with "enter
data" and "exit data" directives I found when adding a new test
(map-subarray-8.f90).
(I'm not sure if this patch fixes all the problems you saw previously
with "alloc", etc. for enter/exit data -- there might still be work to
do there.)
Re-tested with offloading to NVPTX. Does this look OK now?
Thanks,
Julian
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: pointer-and-member-mappings-6.diff --]
[-- Type: text/x-patch, Size: 57191 bytes --]
commit 377dba3ca7dc6fdf86bdb82936e80100ef2dbf4f
Author: Julian Brown <julian@codesourcery.com>
Date: Mon Oct 17 16:44:31 2022 +0000
OpenMP: Pointers and member mappings
This patch changes the mapping node arrangement used for array components
of derived types, e.g.:
type T
integer, pointer, dimension(:) :: arrptr
end type T
type(T) :: tvar
[...]
!$omp target map(tofrom: tvar%arrptr)
This will currently be mapped using three mapping nodes:
GOMP_MAP_TO tvar%arrptr (the descriptor)
GOMP_MAP_TOFROM *tvar%arrptr%data (the actual array data)
GOMP_MAP_ALWAYS_POINTER tvar%arrptr%data (a pointer to the array data)
This follows OMP 5.0, 2.19.7.1 (or OpenMP 5.2, 5.8.3) "map Clause":
"If a list item in a map clause is an associated pointer and the
pointer is not the base pointer of another list item in a map clause
on the same construct, then it is treated as if its pointer target
is implicitly mapped in the same clause. For the purposes of the map
clause, the mapped pointer target is treated as if its base pointer
is the associated pointer."
However, we can also write this:
map(to: tvar%arrptr) map(tofrom: tvar%arrptr(3:8))
and then instead we should follow (OpenMP 5.2, 5.8.3 "map Clause"):
"For map clauses on map-entering constructs, if any list item has a base
pointer for which a corresponding pointer exists in the data environment
upon entry to the region and either a new list item or the corresponding
pointer is created in the device data environment on entry to the region,
then:
1. [Fortran] The corresponding pointer variable is associated with
a pointer target that has the same rank and bounds as the pointer
target of the original pointer, such that the corresponding list item
can be accessed through the pointer in a target region.
2. The corresponding pointer variable becomes an attached pointer
for the corresponding list item."
With this patch you can write the above mappings, and the mapping nodes
used to map pointers to array sections (with descriptors) now look
like this:
1) map(to: tvar%arrptr) -->
GOMP_MAP_TO [implicit] *tvar%arrptr%data (the array data)
GOMP_MAP_TO_PSET tvar%arrptr (the descriptor)
GOMP_MAP_ATTACH_DETACH tvar%arrptr%data
2) map(tofrom: tvar%arrptr(3:8) -->
GOMP_MAP_TOFROM *tvar%arrptr%data(3) (size 8-3+1, etc.)
GOMP_MAP_TO_PSET tvar%arrptr
GOMP_MAP_ATTACH_DETACH tvar%arrptr%data (bias 3, etc.)
In this case, we can determine in the front-end that the
whole-array/pointer mapping (1) is only needed to map the pointer --
so we drop it entirely. (Note also that we set -- early -- the
OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P flag for whole-array-via-pointer
mappings. See below.)
In the middle end, we process mappings using the struct sibling-list
handling machinery by moving the "GOMP_MAP_TO_PSET" node from the middle
of the group of three mapping nodes to the proper sorted position after
the GOMP_MAP_STRUCT mapping:
GOMP_MAP_STRUCT tvar (len: 1)
GOMP_MAP_TO_PSET tvar%arr (size: 64, etc.) <--. moved here
[...] |
GOMP_MAP_TOFROM *tvar%arrptr%data(3) ___|
GOMP_MAP_ATTACH_DETACH tvar%arrptr%data
In another case, if we have an array of derived-type values "dtarr",
and mappings like:
i = 1
j = 1
map(to: dtarr(i)%arrptr) map(tofrom: dtarr(j)%arrptr(3:8))
We still map the same way, but this time we cannot prove that the base
expressions "dtarr(i) and "dtarr(j)" are the same in the front-end.
So we keep both mappings, but we move the "[implicit]" mapping of the
full-array reference to the end of the clause list in gimplify.cc (by
adjusting the topological sorting algorithm):
GOMP_MAP_STRUCT dtvar (len: 2)
GOMP_MAP_TO_PSET dtvar(i)%arrptr
GOMP_MAP_TO_PSET dtvar(j)%arrptr
[...]
GOMP_MAP_TOFROM *dtvar(j)%arrptr%data(3) (size: 8-3+1)
GOMP_MAP_ATTACH_DETACH dtvar(j)%arrptr%data
GOMP_MAP_TO [implicit] *dtvar(i)%arrptr%data(1) (size: whole array)
GOMP_MAP_ATTACH_DETACH dtvar(i)%arrptr%data
Always moving "[implicit]" full-array mappings after array-section
mappings (without that bit set) means that we'll avoid copying the whole
array unnecessarily -- even in cases where we can't prove that the arrays
are the same.
This version of the patch fixes some bugs with "enter data" and "exit
data" directives with this new mapping arrangement. Also now if you
have mappings like this:
#pragma omp target enter data map(to: dv, dv%arr(1:20))
The whole of the derived-type variable "dv" is mapped, so the
GOMP_MAP_TO_PSET for the array-section mapping can be dropped:
GOMP_MAP_TO dv
GOMP_MAP_TO *dv%arr%data
GOMP_MAP_TO_PSET dv%arr <-- deleted (array section mapping)
GOMP_MAP_ATTACH_DETACH dv%arr%data
For struct components, the GOMP_MAP_TO_PSET mapping is turned into
GOMP_MAP_RELEASE at gimplify time for "exit data" directives.
2022-12-15 Julian Brown <julian@codesourcery.com>
gcc/fortran/
* dependency.cc (gfc_omp_expr_prefix_same): New function.
* dependency.h (gfc_omp_expr_prefix_same): Add prototype.
* gfortran.h (gfc_omp_namelist): Add "duplicate_of" field to "u2"
union.
* trans-openmp.cc (dependency.h): Include.
(gfc_trans_omp_array_section): Use GOMP_MAP_TO_PSET unconditionally for
mapping array descriptors.
(gfc_symbol_rooted_namelist): New function.
(gfc_trans_omp_clauses): Check subcomponent and subarray/element
accesses elsewhere in the clause list for pointers to derived types or
array descriptors, and adjust or drop mapping nodes appropriately.
gcc/
* gimplify.cc (omp_tsort_mapping_groups): Process nodes that have
OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P set after those that don't. Add
enter_exit_data parameter.
(omp_resolve_clause_dependencies): Remove GOMP_MAP_TO_PSET mappings if
we're mapping the whole containing derived-type variable.
(omp_accumulate_sibling_list): Adjust GOMP_MAP_TO_PSET handling.
Remove GOMP_MAP_ALWAYS_POINTER handling.
(gimplify_scan_omp_clauses): Pass enter_exit argument to
omp_tsort_mapping_groups. Don't adjust/remove GOMP_MAP_TO_PSET
mappings for derived-type components here.
gcc/testsuite/
* gfortran.dg/gomp/map-subarray-2.f90: New test.
* gfortran.dg/gomp/map-subarray.f90: New test.
libgomp/
* testsuite/libgomp.fortran/map-subarray.f90: New test.
* testsuite/libgomp.fortran/map-subarray-2.f90: New test.
* testsuite/libgomp.fortran/map-subarray-3.f90: New test.
* testsuite/libgomp.fortran/map-subarray-4.f90: New test.
* testsuite/libgomp.fortran/map-subarray-6.f90: New test.
* testsuite/libgomp.fortran/map-subarray-7.f90: New test.
* testsuite/libgomp.fortran/map-subarray-8.f90: New test.
* testsuite/libgomp.fortran/map-subcomponents.f90: New test.
* testsuite/libgomp.fortran/struct-elem-map-1.f90: Adjust for
descriptor-mapping changes. Remove XFAIL.
diff --git a/gcc/fortran/dependency.cc b/gcc/fortran/dependency.cc
index ab3bd36f74ed..1c98f933ff18 100644
--- a/gcc/fortran/dependency.cc
+++ b/gcc/fortran/dependency.cc
@@ -2334,3 +2334,131 @@ gfc_dep_resolver (gfc_ref *lref, gfc_ref *rref, gfc_reverse *reverse,
return fin_dep == GFC_DEP_OVERLAP;
}
+
+/* Check if two refs are equal, for the purposes of checking if one might be
+ the base of the other for OpenMP (target directives). Derived from
+ gfc_dep_resolver. This function is stricter, e.g. indices arr(i) and
+ arr(j) compare as non-equal. */
+
+bool
+gfc_omp_expr_prefix_same (gfc_expr *lexpr, gfc_expr *rexpr)
+{
+ gfc_ref *lref, *rref;
+
+ if (lexpr->symtree && rexpr->symtree)
+ {
+ /* See are_identical_variables above. */
+ if (lexpr->symtree->n.sym->attr.dummy
+ && rexpr->symtree->n.sym->attr.dummy)
+ {
+ /* Dummy arguments: Only check for equal names. */
+ if (lexpr->symtree->n.sym->name != rexpr->symtree->n.sym->name)
+ return false;
+ }
+ else
+ {
+ if (lexpr->symtree->n.sym != rexpr->symtree->n.sym)
+ return false;
+ }
+ }
+ else if (lexpr->base_expr && rexpr->base_expr)
+ {
+ if (gfc_dep_compare_expr (lexpr->base_expr, rexpr->base_expr) != 0)
+ return false;
+ }
+ else
+ return false;
+
+ lref = lexpr->ref;
+ rref = rexpr->ref;
+
+ while (lref && rref)
+ {
+ gfc_dependency fin_dep = GFC_DEP_EQUAL;
+
+ if (lref && lref->type == REF_COMPONENT && lref->u.c.component
+ && strcmp (lref->u.c.component->name, "_data") == 0)
+ lref = lref->next;
+
+ if (rref && rref->type == REF_COMPONENT && rref->u.c.component
+ && strcmp (rref->u.c.component->name, "_data") == 0)
+ rref = rref->next;
+
+ gcc_assert (lref->type == rref->type);
+
+ switch (lref->type)
+ {
+ case REF_COMPONENT:
+ if (lref->u.c.component != rref->u.c.component)
+ return false;
+ break;
+
+ case REF_ARRAY:
+ if (ref_same_as_full_array (lref, rref))
+ break;
+ if (ref_same_as_full_array (rref, lref))
+ break;
+
+ if (lref->u.ar.dimen != rref->u.ar.dimen)
+ {
+ if (lref->u.ar.type == AR_FULL
+ && gfc_full_array_ref_p (rref, NULL))
+ break;
+ if (rref->u.ar.type == AR_FULL
+ && gfc_full_array_ref_p (lref, NULL))
+ break;
+ return false;
+ }
+
+ for (int n = 0; n < lref->u.ar.dimen; n++)
+ {
+ if (lref->u.ar.dimen_type[n] == DIMEN_VECTOR
+ && rref->u.ar.dimen_type[n] == DIMEN_VECTOR
+ && gfc_dep_compare_expr (lref->u.ar.start[n],
+ rref->u.ar.start[n]) == 0)
+ continue;
+ if (lref->u.ar.dimen_type[n] == DIMEN_RANGE
+ && rref->u.ar.dimen_type[n] == DIMEN_RANGE)
+ fin_dep = check_section_vs_section (&lref->u.ar, &rref->u.ar,
+ n);
+ else if (lref->u.ar.dimen_type[n] == DIMEN_ELEMENT
+ && rref->u.ar.dimen_type[n] == DIMEN_RANGE)
+ fin_dep = gfc_check_element_vs_section (lref, rref, n);
+ else if (rref->u.ar.dimen_type[n] == DIMEN_ELEMENT
+ && lref->u.ar.dimen_type[n] == DIMEN_RANGE)
+ fin_dep = gfc_check_element_vs_section (rref, lref, n);
+ else if (lref->u.ar.dimen_type[n] == DIMEN_ELEMENT
+ && rref->u.ar.dimen_type[n] == DIMEN_ELEMENT)
+ {
+ gfc_array_ref l_ar = lref->u.ar;
+ gfc_array_ref r_ar = rref->u.ar;
+ gfc_expr *l_start = l_ar.start[n];
+ gfc_expr *r_start = r_ar.start[n];
+ int i = gfc_dep_compare_expr (r_start, l_start);
+ if (i == 0)
+ fin_dep = GFC_DEP_EQUAL;
+ else
+ return false;
+ }
+ else
+ return false;
+ if (n + 1 < lref->u.ar.dimen
+ && fin_dep != GFC_DEP_EQUAL)
+ return false;
+ }
+
+ if (fin_dep != GFC_DEP_EQUAL
+ && fin_dep != GFC_DEP_OVERLAP)
+ return false;
+
+ break;
+
+ default:
+ gcc_unreachable ();
+ }
+ lref = lref->next;
+ rref = rref->next;
+ }
+
+ return true;
+}
diff --git a/gcc/fortran/dependency.h b/gcc/fortran/dependency.h
index 339be76a8d03..ac94010f84cd 100644
--- a/gcc/fortran/dependency.h
+++ b/gcc/fortran/dependency.h
@@ -40,5 +40,6 @@ int gfc_expr_is_one (gfc_expr *, int);
int gfc_dep_resolver (gfc_ref *, gfc_ref *, gfc_reverse *,
bool identical = false);
int gfc_are_equivalenced_arrays (gfc_expr *, gfc_expr *);
+bool gfc_omp_expr_prefix_same (gfc_expr *, gfc_expr *);
gfc_expr * gfc_discard_nops (gfc_expr *);
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index b541a07e2c7f..612a54f4cc49 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1361,6 +1361,7 @@ typedef struct gfc_omp_namelist
{
struct gfc_omp_namelist_udr *udr;
gfc_namespace *ns;
+ struct gfc_omp_namelist *duplicate_of;
} u2;
struct gfc_omp_namelist *next;
locus where;
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 20a23b2a54cd..c98836131f26 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -40,6 +40,7 @@ along with GCC; see the file COPYING3. If not see
#include "omp-general.h"
#include "omp-low.h"
#include "memmodel.h" /* For MEMMODEL_ enums. */
+#include "dependency.h"
#undef GCC_DIAG_STYLE
#define GCC_DIAG_STYLE __gcc_tdiag__
@@ -2471,28 +2472,15 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
}
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
{
- tree desc_node;
tree type = TREE_TYPE (decl);
ptr2 = gfc_conv_descriptor_data_get (decl);
- desc_node = build_omp_clause (input_location, OMP_CLAUSE_MAP);
- OMP_CLAUSE_DECL (desc_node) = decl;
- OMP_CLAUSE_SIZE (desc_node) = TYPE_SIZE_UNIT (type);
- if (ptr_kind == GOMP_MAP_ALWAYS_POINTER)
- {
- OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO);
- node2 = node;
- node = desc_node; /* Needs to come first. */
- }
- else
- {
- OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO_PSET);
- node2 = desc_node;
- }
- node3 = build_omp_clause (input_location,
- OMP_CLAUSE_MAP);
+ node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
+ OMP_CLAUSE_DECL (node2) = decl;
+ OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+ node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
- OMP_CLAUSE_DECL (node3)
- = gfc_conv_descriptor_data_get (decl);
+ OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (decl);
/* This purposely does not include GOMP_MAP_ALWAYS_POINTER. The extra
cast prevents gimplify.cc from recognising it as being part of the
struct – and adding an 'alloc: for the 'desc.data' pointer, which
@@ -2593,6 +2581,73 @@ handle_iterator (gfc_namespace *ns, stmtblock_t *iter_block, tree block)
return list;
}
+/* To alleviate quadratic behaviour in checking each entry of a
+ gfc_omp_namelist against every other entry, we build a hashtable indexed by
+ gfc_symbol pointer, which we can use in the usual case that a map
+ expression has a symbol as its root term. Return a namelist based on the
+ root symbol used by N, building a new table in SYM_ROOTED_NL using the
+ gfc_omp_namelist N2 (all clauses) if we haven't done so already. */
+
+static gfc_omp_namelist *
+get_symbol_rooted_namelist (hash_map<gfc_symbol *,
+ gfc_omp_namelist *> *&sym_rooted_nl,
+ gfc_omp_namelist *n,
+ gfc_omp_namelist *n2, bool *sym_based)
+{
+ /* Early-out if we have a NULL clause list (e.g. for OpenACC). */
+ if (!n2)
+ return NULL;
+
+ gfc_symbol *use_sym = NULL;
+
+ /* We're only interested in cases where we have an expression, e.g. a
+ component access. */
+ if (n->expr && n->expr->expr_type == EXPR_VARIABLE && n->expr->symtree)
+ use_sym = n->expr->symtree->n.sym;
+
+ *sym_based = false;
+
+ if (!use_sym)
+ return n2;
+
+ if (!sym_rooted_nl)
+ {
+ sym_rooted_nl = new hash_map<gfc_symbol *, gfc_omp_namelist *> ();
+
+ for (; n2 != NULL; n2 = n2->next)
+ {
+ if (!n2->expr
+ || n2->expr->expr_type != EXPR_VARIABLE
+ || !n2->expr->symtree)
+ continue;
+
+ gfc_omp_namelist *nl_copy = gfc_get_omp_namelist ();
+ memcpy (nl_copy, n2, sizeof *nl_copy);
+ nl_copy->u2.duplicate_of = n2;
+ nl_copy->next = NULL;
+
+ gfc_symbol *idx_sym = n2->expr->symtree->n.sym;
+
+ bool existed;
+ gfc_omp_namelist *&entry
+ = sym_rooted_nl->get_or_insert (idx_sym, &existed);
+ if (existed)
+ nl_copy->next = entry;
+ entry = nl_copy;
+ }
+ }
+
+ gfc_omp_namelist **n2_sym = sym_rooted_nl->get (use_sym);
+
+ if (n2_sym)
+ {
+ *sym_based = true;
+ return *n2_sym;
+ }
+
+ return NULL;
+}
+
static tree
gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
locus where, bool declare_simd = false,
@@ -2610,6 +2665,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
if (clauses == NULL)
return NULL_TREE;
+ hash_map<gfc_symbol *, gfc_omp_namelist *> *sym_rooted_nl = NULL;
+
for (list = 0; list < OMP_LIST_NUM; list++)
{
gfc_omp_namelist *n = clauses->lists[list];
@@ -3447,6 +3504,54 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
{
if (pointer || (openacc && allocatable))
{
+ gfc_omp_namelist *n2
+ = openacc ? NULL : clauses->lists[OMP_LIST_MAP];
+
+ bool sym_based;
+ n2 = get_symbol_rooted_namelist (sym_rooted_nl, n,
+ n2, &sym_based);
+
+ /* If the last reference is a pointer to a derived
+ type ("foo%dt_ptr"), check if any subcomponents
+ of the same derived type member are being mapped
+ elsewhere in the clause list ("foo%dt_ptr%x",
+ etc.). If we have such subcomponent mappings,
+ we only create an ALLOC node for the pointer
+ itself, and inhibit mapping the whole derived
+ type. */
+
+ for (; n2 != NULL; n2 = n2->next)
+ {
+ if ((!sym_based && n == n2)
+ || (sym_based && n == n2->u2.duplicate_of)
+ || !n2->expr)
+ continue;
+
+ if (!gfc_omp_expr_prefix_same (n->expr,
+ n2->expr))
+ continue;
+
+ gfc_ref *ref1 = n->expr->ref;
+ gfc_ref *ref2 = n2->expr->ref;
+
+ while (ref1->next && ref2->next)
+ {
+ ref1 = ref1->next;
+ ref2 = ref2->next;
+ }
+
+ if (ref2->next)
+ {
+ inner = build_fold_addr_expr (inner);
+ OMP_CLAUSE_SET_MAP_KIND (node,
+ GOMP_MAP_ALLOC);
+ OMP_CLAUSE_DECL (node) = inner;
+ OMP_CLAUSE_SIZE (node)
+ = TYPE_SIZE_UNIT (TREE_TYPE (inner));
+ goto finalize_map_clause;
+ }
+ }
+
tree data, size;
if (lastref->u.c.component->ts.type == BT_CLASS)
@@ -3494,7 +3599,6 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
{
gomp_map_kind map_kind;
- tree desc_node;
tree type = TREE_TYPE (inner);
tree ptr = gfc_conv_descriptor_data_get (inner);
ptr = build_fold_indirect_ref (ptr);
@@ -3533,21 +3637,69 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
OMP_CLAUSE_SIZE (node)
= fold_build2 (MULT_EXPR, gfc_array_index_type,
OMP_CLAUSE_SIZE (node), elemsz);
- desc_node = build_omp_clause (input_location,
- OMP_CLAUSE_MAP);
- if (openacc)
- OMP_CLAUSE_SET_MAP_KIND (desc_node,
- GOMP_MAP_TO_PSET);
- else
- OMP_CLAUSE_SET_MAP_KIND (desc_node, map_kind);
- OMP_CLAUSE_DECL (desc_node) = inner;
- OMP_CLAUSE_SIZE (desc_node) = TYPE_SIZE_UNIT (type);
- if (openacc)
- node2 = desc_node;
- else
+ node2 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
+ OMP_CLAUSE_DECL (node2) = inner;
+ OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+ if (!openacc)
{
- node2 = node;
- node = desc_node; /* Put first. */
+ gfc_omp_namelist *n2
+ = clauses->lists[OMP_LIST_MAP];
+
+ /* If we don't have a mapping of a smaller part
+ of the array -- or we can't prove that we do
+ statically -- set this flag. If there is a
+ mapping of a smaller part of the array after
+ all, this will turn into a no-op at
+ runtime. */
+ OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (node) = 1;
+
+ bool sym_based;
+ n2 = get_symbol_rooted_namelist (sym_rooted_nl,
+ n, n2,
+ &sym_based);
+
+ bool drop_mapping = false;
+
+ for (; n2 != NULL; n2 = n2->next)
+ {
+ if ((!sym_based && n == n2)
+ || (sym_based && n == n2->u2.duplicate_of)
+ || !n2->expr)
+ continue;
+
+ if (!gfc_omp_expr_prefix_same (n->expr,
+ n2->expr))
+ continue;
+
+ gfc_ref *ref1 = n->expr->ref;
+ gfc_ref *ref2 = n2->expr->ref;
+
+ /* We know ref1 and ref2 overlap. We're
+ interested in whether ref2 describes a
+ smaller part of the array than ref1, which
+ we already know refers to the full
+ array. */
+
+ while (ref1->next && ref2->next)
+ {
+ ref1 = ref1->next;
+ ref2 = ref2->next;
+ }
+
+ if (ref2->next
+ || (ref2->type == REF_ARRAY
+ && (ref2->u.ar.type == AR_ELEMENT
+ || (ref2->u.ar.type
+ == AR_SECTION))))
+ {
+ drop_mapping = true;
+ break;
+ }
+ }
+ if (drop_mapping)
+ continue;
}
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
@@ -3699,6 +3851,23 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
}
}
+ /* Free hashmap if we built it. */
+ if (sym_rooted_nl)
+ {
+ typedef hash_map<gfc_symbol *, gfc_omp_namelist *>::iterator hti;
+ for (hti it = sym_rooted_nl->begin (); it != sym_rooted_nl->end (); ++it)
+ {
+ gfc_omp_namelist *&nl = (*it).second;
+ while (nl)
+ {
+ gfc_omp_namelist *next = nl->next;
+ free (nl);
+ nl = next;
+ }
+ }
+ delete sym_rooted_nl;
+ }
+
if (clauses->if_expr)
{
tree if_var;
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index a8dd298559e8..c59073850382 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9651,18 +9651,45 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
static omp_mapping_group *
omp_tsort_mapping_groups (vec<omp_mapping_group> *groups,
hash_map<tree_operand_hash, omp_mapping_group *>
- *grpmap)
+ *grpmap,
+ bool enter_exit_data)
{
omp_mapping_group *grp, *outlist = NULL, **cursor;
unsigned int i;
+ bool saw_runtime_implicit = false;
cursor = &outlist;
FOR_EACH_VEC_ELT (*groups, i, grp)
{
if (grp->mark != PERMANENT)
- if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp))
- return NULL;
+ {
+ if (OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (*grp->grp_start))
+ {
+ saw_runtime_implicit = true;
+ continue;
+ }
+ if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp))
+ return NULL;
+ }
+ }
+
+ if (!saw_runtime_implicit)
+ return outlist;
+
+ FOR_EACH_VEC_ELT (*groups, i, grp)
+ {
+ if (grp->mark != PERMANENT
+ && OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (*grp->grp_start))
+ {
+ /* Clear the flag for enter/exit data because it is currently
+ meaningless for those operations in libgomp. */
+ if (enter_exit_data)
+ OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (*grp->grp_start) = 0;
+
+ if (!omp_tsort_mapping_groups_1 (&cursor, groups, grpmap, grp))
+ return NULL;
+ }
}
return outlist;
@@ -10060,6 +10087,11 @@ omp_check_mapping_compatibility (location_t loc,
mapping. However, if we have a reference to pointer, make other appropriate
adjustments to the mapping nodes instead.
+ If we have an ATTACH_DETACH node with a Fortran pointer-set (array
+ descriptor) mapping for a derived-type component, and we're also mapping the
+ whole of the derived-type variable on another clause, the pointer-set
+ mapping is removed.
+
If we have a component access but we're also mapping the whole of the
containing struct, drop the former access.
@@ -10239,6 +10271,18 @@ omp_resolve_clause_dependencies (enum tree_code code,
GOMP_MAP_ATTACH_ZLAS for it. */
if (!base_mapped_to && referenced_ptr_node)
OMP_CLAUSE_SET_MAP_KIND (referenced_ptr_node, zlas_kind);
+
+ omp_mapping_group *struct_group;
+ tree desc;
+ if ((desc = OMP_CLAUSE_CHAIN (*grp->grp_start))
+ && OMP_CLAUSE_CODE (desc) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (desc) == GOMP_MAP_TO_PSET
+ && omp_mapped_by_containing_struct (grpmap, decl,
+ &struct_group))
+ /* If we have a pointer set but we're mapping the whole of
+ the containing struct, we can remove the pointer set
+ mapping. */
+ OMP_CLAUSE_CHAIN (*grp->grp_start) = OMP_CLAUSE_CHAIN (desc);
}
else if (TREE_CODE (TREE_TYPE (base_ptr)) == REFERENCE_TYPE
&& (TREE_CODE (TREE_TYPE (TREE_TYPE (base_ptr)))
@@ -10673,11 +10717,19 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
for the purposes of gathering sibling lists, etc. */
/* gcc_assert (base == addr_tokens[base_token]->expr); */
- bool ptr = (OMP_CLAUSE_MAP_KIND (grp_end) == GOMP_MAP_ALWAYS_POINTER);
bool attach_detach = ((OMP_CLAUSE_MAP_KIND (grp_end)
== GOMP_MAP_ATTACH_DETACH)
|| (OMP_CLAUSE_MAP_KIND (grp_end)
== GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION));
+ bool has_descriptor = false;
+ if (OMP_CLAUSE_CHAIN (*grp_start_p) != grp_end)
+ {
+ tree grp_mid = OMP_CLAUSE_CHAIN (*grp_start_p);
+ if (grp_mid
+ && OMP_CLAUSE_CODE (grp_mid) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (grp_mid) == GOMP_MAP_TO_PSET)
+ has_descriptor = true;
+ }
if (!struct_map_to_clause || struct_map_to_clause->get (base) == NULL)
{
@@ -10700,7 +10752,18 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
GOMP_MAP_STRUCT into the middle of the old one. */
tree *insert_node_pos = reprocessing_struct ? *added_tail : grp_start_p;
- if (ptr || attach_detach)
+ if (has_descriptor)
+ {
+ tree desc = OMP_CLAUSE_CHAIN (*grp_start_p);
+ if (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA)
+ OMP_CLAUSE_SET_MAP_KIND (desc, GOMP_MAP_RELEASE);
+ tree sc = *insert_node_pos;
+ OMP_CLAUSE_CHAIN (l) = desc;
+ OMP_CLAUSE_CHAIN (*grp_start_p) = OMP_CLAUSE_CHAIN (desc);
+ OMP_CLAUSE_CHAIN (desc) = sc;
+ *insert_node_pos = l;
+ }
+ else if (attach_detach)
{
tree extra_node;
tree alloc_node
@@ -10931,7 +10994,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
|| OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_ATTACH_DETACH)
sc = &OMP_CLAUSE_CHAIN (*sc);
for (i = 0; i < elems; i++, sc = &OMP_CLAUSE_CHAIN (*sc))
- if ((ptr || attach_detach) && sc == grp_start_p)
+ if (attach_detach && sc == grp_start_p)
break;
else if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF
&& TREE_CODE (OMP_CLAUSE_DECL (*sc)) != INDIRECT_REF
@@ -10987,7 +11050,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
|| (known_eq (coffset, offset)
&& maybe_lt (cbitpos, bitpos)))
{
- if (ptr || attach_detach)
+ if (attach_detach)
scp = sc;
else
break;
@@ -11003,7 +11066,9 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
the list manipulation below. We only need to handle the (pointer
or reference) attach/detach case. */
tree extra_node, alloc_node;
- if (attach_detach)
+ if (has_descriptor)
+ gcc_unreachable ();
+ else if (attach_detach)
alloc_node = build_omp_struct_comp_nodes (code, *grp_start_p,
grp_end, &extra_node);
else
@@ -11036,7 +11101,17 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
return NULL;
}
- if (ptr || attach_detach)
+ if (has_descriptor)
+ {
+ tree desc = OMP_CLAUSE_CHAIN (*grp_start_p);
+ if (code == OMP_TARGET_EXIT_DATA
+ || code == OACC_EXIT_DATA)
+ OMP_CLAUSE_SET_MAP_KIND (desc, GOMP_MAP_RELEASE);
+ omp_siblist_move_node_after (desc,
+ &OMP_CLAUSE_CHAIN (*grp_start_p),
+ scp ? scp : sc);
+ }
+ else if (attach_detach)
{
tree cl = NULL_TREE, extra_node;
tree alloc_node = build_omp_struct_comp_nodes (code, *grp_start_p,
@@ -11457,6 +11532,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
list_p);
omp_mapping_group *outlist = NULL;
+ bool enter_exit = (code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA);
/* Topological sorting may fail if we have duplicate nodes, which
we should have detected and shown an error for already. Skip
@@ -11471,7 +11548,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
groups = omp_gather_mapping_groups (list_p);
grpmap = omp_index_mapping_groups (groups);
- outlist = omp_tsort_mapping_groups (groups, grpmap);
+ outlist = omp_tsort_mapping_groups (groups, grpmap, enter_exit);
outlist = omp_segregate_mapping_groups (outlist);
list_p = omp_reorder_mapping_groups (groups, outlist, list_p);
@@ -11865,27 +11942,45 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
explicit-shape array to a function expecting an assumed-shape
argument. Performing "enter data" inside the called function
would thus be problematic. */
- if (code == OMP_TARGET_EXIT_DATA
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)
- OMP_CLAUSE_SET_MAP_KIND (c, OMP_CLAUSE_MAP_KIND (*prev_list_p)
- == GOMP_MAP_DELETE
- ? GOMP_MAP_DELETE : GOMP_MAP_RELEASE);
- else if ((code == OMP_TARGET_EXIT_DATA || code == OMP_TARGET_UPDATE)
- && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET))
- remove = true;
- else if (code == OMP_TARGET_EXIT_DATA
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
- && OMP_CLAUSE_CHAIN (c)
- && (OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c))
- == OMP_CLAUSE_MAP)
- && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_ATTACH_DETACH)
- || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
- && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL
- (OMP_CLAUSE_CHAIN (c)))) == REFERENCE_TYPE)
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_RELEASE);
+
+ tree desc;
+ if (DECL_P (decl)
+ || (prev_list_p
+ && (desc = OMP_CLAUSE_CHAIN (*prev_list_p))
+ && OMP_CLAUSE_CODE (desc) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (desc) == GOMP_MAP_RELEASE
+ || OMP_CLAUSE_MAP_KIND (desc) == GOMP_MAP_DELETE)
+ && DECL_P (OMP_CLAUSE_DECL (desc))))
+ {
+ /* We only do the transformations here for non-component
+ mappings. Array descriptors for component mappings are
+ handled in omp_accumulate_sibling_list. */
+ if (code == OMP_TARGET_EXIT_DATA
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)
+ {
+ enum gomp_map_kind k
+ = (OMP_CLAUSE_MAP_KIND (*prev_list_p) == GOMP_MAP_DELETE
+ ? GOMP_MAP_DELETE : GOMP_MAP_RELEASE);
+ OMP_CLAUSE_SET_MAP_KIND (c, k);
+ }
+ else if ((code == OMP_TARGET_EXIT_DATA
+ || code == OMP_TARGET_UPDATE)
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET))
+ remove = true;
+ else if (code == OMP_TARGET_EXIT_DATA
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
+ && OMP_CLAUSE_CHAIN (c)
+ && (OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c))
+ == OMP_CLAUSE_MAP)
+ && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_ATTACH_DETACH)
+ || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
+ && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL
+ (OMP_CLAUSE_CHAIN (c)))) == REFERENCE_TYPE)
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_RELEASE);
+ }
if (remove)
break;
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-subarray-2.f90 b/gcc/testsuite/gfortran.dg/gomp/map-subarray-2.f90
new file mode 100644
index 000000000000..033c2538532c
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-subarray-2.f90
@@ -0,0 +1,57 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+type T
+integer, pointer :: arr1(:)
+integer, pointer :: arr2(:)
+integer, pointer :: arr3(:)
+integer, pointer :: arr4(:)
+end type T
+
+type(T) :: tv
+integer, allocatable, target, dimension(:) :: arr
+
+allocate(arr(1:20))
+
+tv%arr1 => arr
+tv%arr2 => arr
+tv%arr3 => arr
+tv%arr4 => arr
+
+!$omp target enter data map(to: tv%arr1)
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target enter data map\(struct:tv \[len: 1\]\) map\(to:tv\.arr1 \[pointer set, len: [0-9]+\]\) map\(to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr1\.data \[bias: 0\]\)} "gimple" } }
+
+!$omp target exit data map(from: tv%arr1)
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target exit data map\(release:tv\.arr1 \[len: [0-9]+\]\) map\(from:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:tv\.arr1\.data \[bias: 0\]\)} "gimple" } }
+
+
+!$omp target enter data map(to: tv%arr2) map(to: tv%arr2(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target enter data map\(struct:tv \[len: 1\]\) map\(to:tv\.arr2 \[pointer set, len: [0-9]+\]\) map\(to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr2\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+!$omp target exit data map(from: tv%arr2) map(from: tv%arr2(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target exit data map\(release:tv\.arr2 \[len: [0-9]+\]\) map\(from:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:tv\.arr2\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+
+!$omp target enter data map(to: tv, tv%arr3(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target enter data map\(to:tv \[len: [0-9]+\]\) map\(to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr3\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+!$omp target exit data map(from: tv, tv%arr3(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target exit data map\(from:tv \[len: [0-9]+\]\) map\(from:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)[_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:tv\.arr3\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+
+!$omp target enter data map(to: tv%arr4(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target enter data map\(struct:tv \[len: 1\]\) map\(to:tv\.arr4 \[pointer set, len: [0-9]+\]\) map\(to:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr4\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+!$omp target exit data map(from: tv%arr4(1:10))
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target exit data map\(release:tv\.arr4 \[len: [0-9]+\]\) map\(from:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(detach:tv\.arr4\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+end
+
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-subarray.f90 b/gcc/testsuite/gfortran.dg/gomp/map-subarray.f90
new file mode 100644
index 000000000000..197888a43365
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-subarray.f90
@@ -0,0 +1,40 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+type T
+integer, pointer :: arr1(:)
+integer, pointer :: arr2(:)
+end type T
+
+type(T) :: tv
+integer, allocatable, target, dimension(:) :: arr
+
+allocate(arr(1:20))
+
+tv%arr1 => arr
+tv%arr2 => arr
+
+!$omp target map(tv%arr1)
+tv%arr1(1) = tv%arr1(1) + 1
+!$omp end target
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target.* map\(struct:tv \[len: 1\]\) map\(to:tv\.arr1 \[pointer set, len: [0-9]+\]\) map\(tofrom:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\[implicit\]\) map\(attach:tv\.arr1\.data \[bias: 0\]\)} "gimple" } }
+
+!$omp target map(tv%arr2) map(tv%arr2(1:10))
+tv%arr2(1) = tv%arr2(1) + 1
+!$omp end target
+
+!$omp target map(tv%arr2(1:10))
+tv%arr2(1) = tv%arr2(1) + 1
+!$omp end target
+
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target.* map\(struct:tv \[len: 1\]\) map\(to:tv\.arr2 \[pointer set, len: [0-9]+\]\) map\(tofrom:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr2\.data \[bias: [^\]]+\]\)} 2 "gimple" } }
+
+!$omp target map(tv, tv%arr2(1:10))
+tv%arr2(1) = tv%arr2(1) + 1
+!$omp end target
+
+! { dg-final { scan-tree-dump {(?n)#pragma omp target.* map\(tofrom:tv \[len: [0-9]+\]\) map\(tofrom:MEM <integer\(kind=4\)\[0:\]> \[\(integer\(kind=4\)\[0:\] \*\)_[0-9]+\] \[len: _[0-9]+\]\) map\(attach:tv\.arr2\.data \[bias: [^\]]+\]\)} "gimple" } }
+
+end
+
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-2.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-2.f90
new file mode 100644
index 000000000000..02f08c52a8c3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-2.f90
@@ -0,0 +1,108 @@
+! { dg-do run }
+
+program myprog
+type u
+ integer, dimension (:), pointer :: tarr1
+ integer, dimension (:), pointer :: tarr2
+ integer, dimension (:), pointer :: tarr3
+end type u
+
+type(u) :: myu1, myu2, myu3
+
+integer, dimension (12), target :: myarray1
+integer, dimension (12), target :: myarray2
+integer, dimension (12), target :: myarray3
+integer, dimension (12), target :: myarray4
+integer, dimension (12), target :: myarray5
+integer, dimension (12), target :: myarray6
+integer, dimension (12), target :: myarray7
+integer, dimension (12), target :: myarray8
+integer, dimension (12), target :: myarray9
+
+myu1%tarr1 => myarray1
+myu1%tarr2 => myarray2
+myu1%tarr3 => myarray3
+myu2%tarr1 => myarray4
+myu2%tarr2 => myarray5
+myu2%tarr3 => myarray6
+myu3%tarr1 => myarray7
+myu3%tarr2 => myarray8
+myu3%tarr3 => myarray9
+
+myu1%tarr1 = 0
+myu1%tarr2 = 0
+myu1%tarr3 = 0
+myu2%tarr1 = 0
+myu2%tarr2 = 0
+myu2%tarr3 = 0
+myu3%tarr1 = 0
+myu3%tarr2 = 0
+myu3%tarr3 = 0
+
+!$omp target map(to:myu1%tarr1) map(tofrom:myu1%tarr1(:)) &
+!$omp& map(to:myu1%tarr2) map(tofrom:myu1%tarr2(:)) &
+!$omp& map(to:myu1%tarr3) map(tofrom:myu1%tarr3(:)) &
+!$omp& map(to:myu2%tarr1) map(tofrom:myu2%tarr1(:)) &
+!$omp& map(to:myu2%tarr2) map(tofrom:myu2%tarr2(:)) &
+!$omp& map(to:myu2%tarr3) map(tofrom:myu2%tarr3(:)) &
+!$omp& map(to:myu3%tarr1) map(tofrom:myu3%tarr1(:)) &
+!$omp& map(to:myu3%tarr2) map(tofrom:myu3%tarr2(:)) &
+!$omp& map(to:myu3%tarr3) map(tofrom:myu3%tarr3(:))
+myu1%tarr1(1) = myu1%tarr1(1) + 1
+myu2%tarr1(1) = myu2%tarr1(1) + 1
+myu3%tarr1(1) = myu3%tarr1(1) + 1
+!$omp end target
+
+!$omp target map(to:myu1%tarr1) map(tofrom:myu1%tarr1(1:2)) &
+!$omp& map(to:myu1%tarr2) map(tofrom:myu1%tarr2(1:2)) &
+!$omp& map(to:myu1%tarr3) map(tofrom:myu1%tarr3(1:2)) &
+!$omp& map(to:myu2%tarr1) map(tofrom:myu2%tarr1(1:2)) &
+!$omp& map(to:myu2%tarr2) map(tofrom:myu2%tarr2(1:2)) &
+!$omp& map(to:myu2%tarr3) map(tofrom:myu2%tarr3(1:2)) &
+!$omp& map(to:myu3%tarr1) map(tofrom:myu3%tarr1(1:2)) &
+!$omp& map(to:myu3%tarr2) map(tofrom:myu3%tarr2(1:2)) &
+!$omp& map(to:myu3%tarr3) map(tofrom:myu3%tarr3(1:2))
+myu1%tarr2(1) = myu1%tarr2(1) + 1
+myu2%tarr2(1) = myu2%tarr2(1) + 1
+myu3%tarr2(1) = myu3%tarr2(1) + 1
+!$omp end target
+
+!$omp target map(to:myu1%tarr1) map(tofrom:myu1%tarr1(1)) &
+!$omp& map(to:myu1%tarr2) map(tofrom:myu1%tarr2(1)) &
+!$omp& map(to:myu1%tarr3) map(tofrom:myu1%tarr3(1)) &
+!$omp& map(to:myu2%tarr1) map(tofrom:myu2%tarr1(1)) &
+!$omp& map(to:myu2%tarr2) map(tofrom:myu2%tarr2(1)) &
+!$omp& map(to:myu2%tarr3) map(tofrom:myu2%tarr3(1)) &
+!$omp& map(to:myu3%tarr1) map(tofrom:myu3%tarr1(1)) &
+!$omp& map(to:myu3%tarr2) map(tofrom:myu3%tarr2(1)) &
+!$omp& map(to:myu3%tarr3) map(tofrom:myu3%tarr3(1))
+myu1%tarr3(1) = myu1%tarr3(1) + 1
+myu2%tarr3(1) = myu2%tarr3(1) + 1
+myu3%tarr3(1) = myu3%tarr3(1) + 1
+!$omp end target
+
+!$omp target map(tofrom:myu1%tarr1) &
+!$omp& map(tofrom:myu1%tarr2) &
+!$omp& map(tofrom:myu1%tarr3) &
+!$omp& map(tofrom:myu2%tarr1) &
+!$omp& map(tofrom:myu2%tarr2) &
+!$omp& map(tofrom:myu2%tarr3) &
+!$omp& map(tofrom:myu3%tarr1) &
+!$omp& map(tofrom:myu3%tarr2) &
+!$omp& map(tofrom:myu3%tarr3)
+myu1%tarr2(1) = myu1%tarr2(1) + 1
+myu2%tarr2(1) = myu2%tarr2(1) + 1
+myu3%tarr2(1) = myu3%tarr2(1) + 1
+!$omp end target
+
+if (myu1%tarr1(1).ne.1) stop 1
+if (myu2%tarr1(1).ne.1) stop 2
+if (myu3%tarr1(1).ne.1) stop 3
+if (myu1%tarr2(1).ne.2) stop 4
+if (myu2%tarr2(1).ne.2) stop 5
+if (myu3%tarr2(1).ne.2) stop 6
+if (myu1%tarr3(1).ne.1) stop 7
+if (myu2%tarr3(1).ne.1) stop 8
+if (myu3%tarr3(1).ne.1) stop 9
+
+end program myprog
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-3.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-3.f90
new file mode 100644
index 000000000000..318e77ea44ff
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-3.f90
@@ -0,0 +1,62 @@
+! { dg-do run }
+
+module mymod
+type G
+integer :: x, y
+integer, pointer :: arr(:)
+integer :: z
+end type G
+end module mymod
+
+program myprog
+use mymod
+
+integer, target :: arr1(10)
+integer, target :: arr2(10)
+integer, target :: arr3(10)
+type(G), dimension(3) :: gvar
+
+integer :: i, j
+
+gvar(1)%arr => arr1
+gvar(2)%arr => arr2
+gvar(3)%arr => arr3
+
+gvar(1)%arr = 0
+gvar(2)%arr = 0
+gvar(3)%arr = 0
+
+i = 1
+j = 1
+
+! Here 'gvar(i)' and 'gvar(j)' are the same element, so this should work.
+! This generates a whole-array mapping for gvar(i)%arr, but with the
+! "runtime implicit" bit set so the smaller subarray gvar(j)%arr(1:5) takes
+! precedence.
+
+!$omp target map(gvar(i)%arr, gvar(j)%arr(1:5))
+gvar(i)%arr(1) = gvar(i)%arr(1) + 1
+gvar(j)%arr(1) = gvar(j)%arr(1) + 2
+!$omp end target
+
+!$omp target map(gvar(i)%arr(1:5), gvar(j)%arr)
+gvar(i)%arr(1) = gvar(i)%arr(1) + 3
+gvar(j)%arr(1) = gvar(j)%arr(1) + 4
+!$omp end target
+
+! For these ones, we know the array index is the same, so we can just
+! drop the whole-array mapping.
+
+!$omp target map(gvar(i)%arr, gvar(i)%arr(1:5))
+gvar(i)%arr(1) = gvar(i)%arr(1) + 1
+gvar(i)%arr(1) = gvar(j)%arr(1) + 2
+!$omp end target
+
+!$omp target map(gvar(i)%arr(1:5), gvar(i)%arr)
+gvar(i)%arr(1) = gvar(i)%arr(1) + 3
+gvar(i)%arr(1) = gvar(j)%arr(1) + 4
+!$omp end target
+
+if (gvar(1)%arr(1).ne.20) stop 1
+
+end program myprog
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-4.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-4.f90
new file mode 100644
index 000000000000..5d15808f0da7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-4.f90
@@ -0,0 +1,35 @@
+! { dg-do run }
+
+type t
+ integer, pointer :: p(:)
+end type t
+
+type(t) :: var(2)
+
+allocate (var(1)%p, source=[1,2,3,5])
+allocate (var(2)%p, source=[2,3,5])
+
+!$omp target map(var(1)%p, var(2)%p)
+var(1)%p(1) = 5
+var(2)%p(2) = 7
+!$omp end target
+
+!$omp target map(var(1)%p(1:3), var(1)%p, var(2)%p)
+var(1)%p(1) = var(1)%p(1) + 1
+var(2)%p(2) = var(2)%p(2) + 1
+!$omp end target
+
+!$omp target map(var(1)%p, var(2)%p, var(2)%p(1:3))
+var(1)%p(1) = var(1)%p(1) + 1
+var(2)%p(2) = var(2)%p(2) + 1
+!$omp end target
+
+!$omp target map(var(1)%p, var(1)%p(1:3), var(2)%p, var(2)%p(2))
+var(1)%p(1) = var(1)%p(1) + 1
+var(2)%p(2) = var(2)%p(2) + 1
+!$omp end target
+
+if (var(1)%p(1).ne.8) stop 1
+if (var(2)%p(2).ne.10) stop 2
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-6.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-6.f90
new file mode 100644
index 000000000000..9f0edf70890e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-6.f90
@@ -0,0 +1,26 @@
+! { dg-do run }
+
+type t
+ integer, pointer :: p(:)
+ integer, pointer :: p2(:)
+end type t
+
+type(t) :: var
+integer, target :: tgt(5), tgt2(1000)
+var%p => tgt
+var%p2 => tgt2
+
+p = 0
+p2 = 0
+
+!$omp target map(tgt, tgt2(4:6), var)
+ var%p(1) = 5
+ var%p2(5) = 7
+!$omp end target
+
+if (var%p(1).ne.5) stop 1
+if (var%p2(5).ne.7) stop 2
+
+end
+
+! { dg-shouldfail "" { offload_device_nonshared_as } }
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-7.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-7.f90
new file mode 100644
index 000000000000..42da72961069
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-7.f90
@@ -0,0 +1,29 @@
+type t
+integer, pointer :: p2(:)
+end type t
+
+integer, target :: A(5)
+integer, pointer :: p(:), p2(:)
+type(t) :: var
+
+allocate(p2(1:20))
+p => A
+var%p2 => p2
+
+A = 0
+p2 = 0
+
+! These arrays "share original storage", so are unsupported. This will
+! (correctly) fail with a non-shared address space.
+
+!$omp target map(A(3:4), p2(4:8), p, var%p2)
+A(3) = A(3) + 1
+p2(4) = p2(4) + 2
+!$omp end target
+
+if (A(3).ne.1) stop 1
+if (p2(4).ne.2) stop 2
+
+end program
+
+! { dg-shouldfail "" { offload_device_nonshared_as } }
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-8.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-8.f90
new file mode 100644
index 000000000000..a47360e10ec3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray-8.f90
@@ -0,0 +1,47 @@
+! { dg-do run }
+
+type F
+integer, pointer :: mem(:)
+end type F
+
+type(F) :: fv
+integer, allocatable, target :: arr(:)
+
+allocate(arr(1:20))
+
+fv%mem => arr
+fv%mem = 0
+
+!$omp target enter data map(to: fv%mem(1:10))
+!$omp target map(alloc: fv%mem)
+fv%mem(1) = fv%mem(1) + 1
+!$omp end target
+!$omp target exit data map(from: fv%mem(1:10))
+
+if (fv%mem(1).ne.1) stop 1
+
+!$omp target enter data map(to: fv, fv%mem(1:10))
+!$omp target
+fv%mem(1) = fv%mem(1) + 1
+!$omp end target
+!$omp target exit data map(from: fv, fv%mem(1:10))
+
+if (fv%mem(1).ne.2) stop 2
+
+!$omp target enter data map(to: fv%mem, fv%mem(1:10))
+!$omp target
+fv%mem(1) = fv%mem(1) + 1
+!$omp end target
+!$omp target exit data map(from: fv%mem, fv%mem(1:10))
+
+if (fv%mem(1).ne.3) stop 3
+
+!$omp target enter data map(to: fv%mem)
+!$omp target
+fv%mem(1) = fv%mem(1) + 1
+!$omp end target
+!$omp target exit data map(from: fv%mem)
+
+if (fv%mem(1).ne.4) stop 4
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray.f90
new file mode 100644
index 000000000000..85f5af3a2a6c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subarray.f90
@@ -0,0 +1,33 @@
+! { dg-do run }
+
+program myprog
+type u
+ integer, dimension (:), pointer :: tarr
+end type u
+
+type(u) :: myu
+integer, dimension (12), target :: myarray
+
+myu%tarr => myarray
+
+myu%tarr = 0
+
+!$omp target map(to:myu%tarr) map(tofrom:myu%tarr(:))
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+!$omp target map(to:myu%tarr) map(tofrom:myu%tarr(1:2))
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+!$omp target map(to:myu%tarr) map(tofrom:myu%tarr(1))
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+!$omp target map(tofrom:myu%tarr)
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+if (myu%tarr(1).ne.4) stop 1
+
+end program myprog
diff --git a/libgomp/testsuite/libgomp.fortran/map-subcomponents.f90 b/libgomp/testsuite/libgomp.fortran/map-subcomponents.f90
new file mode 100644
index 000000000000..c7f90131cbae
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/map-subcomponents.f90
@@ -0,0 +1,32 @@
+! { dg-do run }
+
+module mymod
+type F
+integer :: a, b, c
+integer, dimension(10) :: d
+end type F
+
+type G
+integer :: x, y
+type(F), pointer :: myf
+integer :: z
+end type G
+end module mymod
+
+program myprog
+use mymod
+
+type(F), target :: ftmp
+type(G) :: gvar
+
+gvar%myf => ftmp
+
+gvar%myf%d = 0
+
+!$omp target map(to:gvar%myf) map(tofrom: gvar%myf%b, gvar%myf%d)
+gvar%myf%d(1) = gvar%myf%d(1) + 1
+!$omp end target
+
+if (gvar%myf%d(1).ne.1) stop 1
+
+end program myprog
diff --git a/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90 b/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90
index 7f3d8174f97b..b1d696656c0f 100644
--- a/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90
@@ -36,6 +36,10 @@ program main
call six ()
call seven ()
call eight ()
+ call nine ()
+ call ten ()
+ call eleven ()
+ call twelve ()
contains
! Implicitly mapped – but no pointers are mapped
@@ -408,7 +412,180 @@ contains
!$omp end target
end subroutine eight
-end program main
+ ! This is "subroutine four" but with explicit base-pointer mappings
+ ! (var%f, etc.).
+ subroutine nine()
+ type(t2) :: var
-! Fixed by the "Fortran pointers and member mappings" patch
-! { dg-xfail-run-if TODO { offload_device_nonshared_as } }
+ print '(g0)', '==== TESTCASE "nine" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
+ uni1 = 4_"abcde", &
+ uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str4, source=["Let's", "Go!!!"])
+ allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
+
+! !$omp target map(tofrom: var%d(4:7), var%f(2:3), var%str2(2:3)) &
+! !$omp& map(tofrom: var%str4(2:2), var%uni2(2:3), var%uni4(2:2))
+ !$omp target map(to: var%f) map(tofrom: var%d(4:7), var%f(2:3), &
+ !$omp& var%str2(2:3), var%uni2(2:3))
+ if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
+ if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+
+ if (.not. associated (var%f)) stop 9
+ if (size (var%f) /= 4) stop 10
+ if (any (var%f(2:3) /= [33, 44])) stop 11
+! if (.not. associated (var%str4)) stop 15
+! if (len (var%str4) /= 5) stop 16
+! if (size (var%str4) /= 2) stop 17
+! if (var%str4(2) /= "Go!!!") stop 18
+
+ if (any (var%uni2(2:3) /= [4_"67890", 4_"ABCDE"])) stop 19
+! if (.not. associated (var%uni4)) stop 20
+! if (len (var%uni4) /= 5) stop 21
+! if (size (var%uni4) /= 2) stop 22
+! if (var%uni4(2) /= "Go!!!") stop 23
+ !$omp end target
+
+ deallocate(var%f, var%str4)
+ end subroutine nine
+
+ ! This is "subroutine five" but with explicit base-pointer mappings.
+ subroutine ten()
+ type(t2) :: var
+
+ print '(g0)', '==== TESTCASE "ten" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
+ uni1 = 4_"abcde", &
+ uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str4, source=["Let's", "Go!!!"])
+
+ !$omp target map(tofrom: var%d(4:7))
+ if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
+ !$omp end target
+ !$omp target map(tofrom: var%str2(2:3))
+ if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+ !$omp end target
+
+ !$omp target map(to: var%f) map(tofrom: var%f(2:3))
+ if (.not. associated (var%f)) stop 9
+ if (size (var%f) /= 4) stop 10
+ if (any (var%f(2:3) /= [33, 44])) stop 11
+ !$omp end target
+! !$omp target map(tofrom: var%str4(2:2))
+! if (.not. associated (var%str4)) stop 15
+! if (len (var%str4) /= 5) stop 16
+! if (size (var%str4) /= 2) stop 17
+! if (var%str4(2) /= "Go!!!") stop 18
+! !$omp end target
+! !$omp target map(tofrom: var%uni4(2:2))
+! if (.not. associated (var%uni4)) stop 15
+! if (len (var%uni4) /= 5) stop 16
+! if (size (var%uni4) /= 2) stop 17
+! if (var%uni4(2) /= 4_"Go!!!") stop 18
+! !$omp end target
+
+ deallocate(var%f, var%str4)
+ end subroutine ten
+
+ ! This is "subroutine six" but with explicit base pointer mappings.
+ subroutine eleven()
+ type(t2) :: var
+
+ print '(g0)', '==== TESTCASE "eleven" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
+ uni1 = 4_"abcde", &
+ uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str4, source=["Let's", "Go!!!"])
+ allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
+
+! !$omp target map(tofrom: var%d(5), var%f(3), var%str2(3), &
+! !$omp var%str4(2), var%uni2(3), var%uni4(2))
+ !$omp target map(to: var%f) map(tofrom: var%d(5), var%f(3), &
+ !$omp& var%str2(3), var%uni2(3))
+ if (var%d(5) /= -3*5) stop 4
+ if (var%str2(3) /= "ABCDE") stop 6
+ if (var%uni2(3) /= 4_"ABCDE") stop 7
+
+ if (.not. associated (var%f)) stop 9
+ if (size (var%f) /= 4) stop 10
+ if (var%f(3) /= 44) stop 11
+! if (.not. associated (var%str4)) stop 15
+! if (len (var%str4) /= 5) stop 16
+! if (size (var%str4) /= 2) stop 17
+! if (var%str4(2) /= "Go!!!") stop 18
+! if (.not. associated (var%uni4)) stop 19
+! if (len (var%uni4) /= 5) stop 20
+! if (size (var%uni4) /= 2) stop 21
+! if (var%uni4(2) /= 4_"Go!!!") stop 22
+ !$omp end target
+
+ deallocate(var%f, var%str4, var%uni4)
+ end subroutine eleven
+
+ ! This is "subroutine seven" but with explicit base-pointer mappings.
+ subroutine twelve()
+ type(t2) :: var
+
+ print '(g0)', '==== TESTCASE "twelve" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"], &
+ uni1 = 4_"abcde", &
+ uni2 = [4_"12345", 4_"67890", 4_"ABCDE", 4_"FGHIJ"])
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str4, source=["Let's", "Go!!!"])
+ allocate (var%uni4, source=[4_"Let's", 4_"Go!!!"])
+
+ !$omp target map(tofrom: var%d(5))
+ if (var%d(5) /= (-3*5)) stop 4
+ !$omp end target
+ !$omp target map(tofrom: var%str2(2:3))
+ if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+ !$omp end target
+ !$omp target map(tofrom: var%uni2(2:3))
+ if (any (var%uni2(2:3) /= [4_"67890", 4_"ABCDE"])) stop 7
+ !$omp end target
+
+ !$omp target map(to: var%f) map(tofrom: var%f(2:3))
+ if (.not. associated (var%f)) stop 9
+ if (size (var%f) /= 4) stop 10
+ if (any (var%f(2:3) /= [33, 44])) stop 11
+ !$omp end target
+! !$omp target map(tofrom: var%str4(2:2))
+! if (.not. associated (var%str4)) stop 15
+! if (len (var%str4) /= 5) stop 16
+! if (size (var%str4) /= 2) stop 17
+! if (var%str4(2) /= "Go!!!") stop 18
+! !$omp end target
+! !$omp target map(tofrom: var%uni4(2:2))
+! if (.not. associated (var%uni4)) stop 15
+! if (len (var%uni4) /= 5) stop 16
+! if (size (var%uni4) /= 2) stop 17
+! if (var%uni4(2) /= 4_"Go!!!") stop 18
+! !$omp end target
+
+ deallocate(var%f, var%str4, var%uni4)
+ end subroutine twelve
+
+end program main
next prev parent reply other threads:[~2022-12-15 14:55 UTC|newest]
Thread overview: 14+ messages / expand[flat|nested] mbox.gz Atom feed top
2022-10-18 10:39 [PATCH v5 0/4] OpenMP/OpenACC: Fortran array descriptor mappings Julian Brown
2022-10-18 10:39 ` [PATCH v5 1/4] OpenMP/OpenACC: Reindent TO/FROM/_CACHE_ stanza in {c_}finish_omp_clause Julian Brown
2022-10-18 10:39 ` [PATCH v5 2/4] OpenMP/OpenACC: Rework clause expansion and nested struct handling Julian Brown
2022-10-18 12:33 ` Julian Brown
2022-10-18 12:33 ` Julian Brown
2022-12-07 14:54 ` Tobias Burnus
2022-12-07 15:16 ` Julian Brown
2022-12-07 16:13 ` Tobias Burnus
2022-12-12 15:19 ` Julian Brown
2022-10-18 10:39 ` [PATCH v5 3/4] OpenMP: Pointers and member mappings Julian Brown
2022-12-07 16:31 ` Tobias Burnus
2022-12-15 14:54 ` Julian Brown [this message]
2022-12-15 16:46 ` Julian Brown
2022-10-18 10:39 ` [PATCH v5 4/4] OpenMP/OpenACC: Unordered/non-constant component offset runtime diagnostic Julian Brown
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=20221215145458.458074ca@squid.athome \
--to=julian@codesourcery.com \
--cc=fortran@gcc.gnu.org \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
--cc=tobias@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).