From: Kwok Cheung Yeung <[email protected]>
This adds support for iterators in map clauses within OpenMP
'target' constructs in Fortran.
Some special handling for struct field maps has been added to libgomp in
order to handle arrays of derived types.
gcc/fortran/
* dump-parse-tree.cc (show_omp_namelist): Add iterator support for
OMP_LIST_MAP.
* match.cc (gfc_free_omp_namelist): Free namespace in namelist for
OMP_LIST_MAP.
* openmp.cc (gfc_match_omp_clauses): Parse 'iterator' modifier for
'map' clause.
(resolve_omp_clauses): Resolve iterators for OMP_LIST_MAP.
* trans-openmp.cc: Include tree-ssa-loop-niter.h.
(gfc_trans_omp_array_section): Add iterator argument. Replace
instances of iterator variables with the initial value when
computing biases.
(gfc_trans_omp_clauses): Handle iterators in OMP_LIST_MAP clauses.
Add expressions to iter_block rather than block. Do not apply
iterators to firstprivate maps. Pass iterator to
gfc_trans_omp_array_section.
gcc/
* gimplify.cc (compute_omp_iterator_count): Account for difference
in loop boundaries in Fortran.
(build_omp_iterator_loop): Change upper boundary condition for
Fortran. Insert block statements into innermost loop.
(remove_unused_omp_iterator_vars): Copy block subblocks of old
iterator to new iterator and remove original.
(extract_base_bit_offset): Add iterator argument. Remove iterator
variables from base.
(omp_accumulate_sibling_list): Add iterator argument to
extract_base_bit_offset.
* tree-pretty-print.cc (dump_block_node): Ignore BLOCK_SUBBLOCKS
containing iterator block statements.
gcc/testsuite/
* gfortran.dg/gomp/target-map-iterators-1.f90: New.
* gfortran.dg/gomp/target-map-iterators-2.f90: New.
* gfortran.dg/gomp/target-map-iterators-3.f90: New.
* gfortran.dg/gomp/target-map-iterators-4.f90: New.
* gfortran.dg/gomp/target-map-iterators-6.f90: New.
libgomp/
* target.c (kind_to_name): Handle GOMP_MAP_STRUCT and
GOMP_MAP_STRUCT_UNORD.
(gomp_add_map): New.
(gomp_merge_iterator_maps): Expand fields of a struct mapping
breadth-first.
* testsuite/libgomp.fortran/target-map-iterators-1.f90: New.
* testsuite/libgomp.fortran/target-map-iterators-2.f90: New.
* testsuite/libgomp.fortran/target-map-iterators-3.f90: New.
Co-authored-by: Andrew Stubbs <[email protected]>
Co-authored-by: Sandra Loosemore <[email protected]>
---
gcc/fortran/dump-parse-tree.cc | 9 +-
gcc/fortran/match.cc | 3 +-
gcc/fortran/openmp.cc | 30 ++++-
gcc/fortran/trans-openmp.cc | 108 ++++++++++++++----
gcc/gimplify.cc | 70 +++++++++---
.../gomp/target-map-iterators-1.f90 | 26 +++++
.../gomp/target-map-iterators-2.f90 | 30 +++++
.../gomp/target-map-iterators-3.f90 | 24 ++++
.../gomp/target-map-iterators-4.f90 | 28 +++++
.../gomp/target-map-iterators-6.f90 | 26 +++++
gcc/tree-pretty-print.cc | 4 +-
libgomp/target.c | 82 +++++++++----
.../target-map-iterators-1.f90 | 45 ++++++++
.../target-map-iterators-2.f90 | 45 ++++++++
.../target-map-iterators-3.f90 | 50 ++++++++
15 files changed, 513 insertions(+), 67 deletions(-)
create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-map-iterators-1.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-map-iterators-2.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-map-iterators-4.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-map-iterators-6.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/target-map-iterators-1.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/target-map-iterators-2.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/target-map-iterators-3.f90
diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
index e7bfa480a2c..03e1ec0f78c 100644
--- a/gcc/fortran/dump-parse-tree.cc
+++ b/gcc/fortran/dump-parse-tree.cc
@@ -1485,7 +1485,8 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n)
for (; n; n = n->next)
{
gfc_current_ns = ns_curr;
- if (list_type == OMP_LIST_AFFINITY || list_type == OMP_LIST_DEPEND)
+ if (list_type == OMP_LIST_AFFINITY || list_type == OMP_LIST_DEPEND
+ || list_type == OMP_LIST_MAP)
{
gfc_current_ns = n->u2.ns ? n->u2.ns : ns_curr;
if (n->u2.ns != ns_iter)
@@ -1497,8 +1498,12 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n)
fputs ("AFFINITY (", dumpfile);
else if (n->u.depend_doacross_op == OMP_DOACROSS_SINK_FIRST)
fputs ("DOACROSS (", dumpfile);
- else
+ else if (list_type == OMP_LIST_DEPEND)
fputs ("DEPEND (", dumpfile);
+ else if (list_type == OMP_LIST_MAP)
+ fputs ("MAP (", dumpfile);
+ else
+ gcc_unreachable ();
}
if (n->u2.ns)
{
diff --git a/gcc/fortran/match.cc b/gcc/fortran/match.cc
index d892a4588b2..6cb2fc06b3c 100644
--- a/gcc/fortran/match.cc
+++ b/gcc/fortran/match.cc
@@ -6341,7 +6341,8 @@ gfc_free_namelist (gfc_namelist *name)
void
gfc_free_omp_namelist (gfc_omp_namelist *name, enum gfc_omp_list_type list)
{
- bool free_ns = (list == OMP_LIST_AFFINITY || list == OMP_LIST_DEPEND);
+ bool free_ns = (list == OMP_LIST_AFFINITY || list == OMP_LIST_DEPEND
+ || list == OMP_LIST_MAP);
bool free_align_allocator = (list == OMP_LIST_ALLOCATE);
bool free_mem_traits_space = (list == OMP_LIST_USES_ALLOCATORS);
bool free_init = (list == OMP_LIST_INIT);
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 5c3c373c27d..a3adae8ae0d 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -3735,11 +3735,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const
omp_mask mask,
int close_modifier = 0;
int present_modifier = 0;
int mapper_modifier = 0;
+ int iterator_modifier = 0;
+ gfc_namespace *ns_iter = NULL, *ns_curr = gfc_current_ns;
locus second_always_locus = old_loc2;
locus second_close_locus = old_loc2;
locus second_mapper_locus = old_loc2;
locus second_present_locus = old_loc2;
char mapper_id[GFC_MAX_SYMBOL_LEN + 1] = { '\0' };
+ locus second_iterator_locus = old_loc2;
for (;;)
{
@@ -3767,6 +3770,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const
omp_mask mask,
if (m != MATCH_YES)
goto error;
}
+ else if (gfc_match_iterator (&ns_iter, true) == MATCH_YES)
+ {
+ if (iterator_modifier++ == 1)
+ second_iterator_locus = current_locus;
+ }
else
break;
if (gfc_match (", ") != MATCH_YES)
@@ -3835,17 +3843,25 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const
omp_mask mask,
&second_mapper_locus);
break;
}
+ if (iterator_modifier > 1)
+ {
+ gfc_error ("too many %<iterator%> modifiers at %L",
+ &second_iterator_locus);
+ break;
+ }
head = NULL;
- if (gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_MAP],
- false, NULL, &head,
- true, true) == MATCH_YES)
+ if (ns_iter)
+ gfc_current_ns = ns_iter;
+ m = gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_MAP],
+ false, NULL, &head, true, true);
+ gfc_current_ns = ns_curr;
+ if (m == MATCH_YES)
{
gfc_omp_namelist *n;
for (n = *head; n; n = n->next)
{
n->u.map.op = map_op;
-
gfc_typespec *ts;
if (n->expr)
ts = &n->expr->ts;
@@ -3858,6 +3874,9 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const
omp_mask mask,
{
n->u2.udm.udm = udm;
}
+ n->u2.ns = ns_iter;
+ if (ns_iter)
+ ns_iter->refs++;
}
continue;
}
@@ -10014,7 +10033,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses
*omp_clauses,
case OMP_LIST_CACHE:
for (; n != NULL; n = n->next)
{
- if ((list == OMP_LIST_DEPEND || list == OMP_LIST_AFFINITY)
+ if ((list == OMP_LIST_DEPEND || list == OMP_LIST_AFFINITY
+ || list == OMP_LIST_MAP)
&& n->u2.ns && !n->u2.ns->resolved)
{
n->u2.ns->resolved = 1;
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 022f14b6954..c3ff30092c6 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -27,6 +27,7 @@ along with GCC; see the file COPYING3. If not see
#include "gfortran.h"
#include "basic-block.h"
#include "tree-ssa.h"
+#include "tree-ssa-loop-niter.h" /* for simplify_replace_tree. */
#include "function.h"
#include "gimple.h"
#include "gimple-expr.h"
@@ -3339,7 +3340,8 @@ static void
gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op,
gfc_omp_namelist *n, tree decl, bool element,
bool openmp, gomp_map_kind ptr_kind, tree &node,
- tree &node2, tree &node3, tree &node4)
+ tree &node2, tree &node3, tree &node4,
+ tree iterator)
{
gfc_se se;
tree ptr, ptr2;
@@ -3494,6 +3496,13 @@ gfc_trans_omp_array_section (stmtblock_t *block,
gfc_exec_op op,
OMP_CLAUSE_DECL (node3) = decl;
}
ptr2 = fold_convert (ptrdiff_type_node, ptr2);
+ for (tree it = iterator; it; it = TREE_CHAIN (it))
+ {
+ ptr = simplify_replace_tree (ptr, TREE_VEC_ELT (it, 0),
+ TREE_VEC_ELT (it, 1));
+ ptr2 = simplify_replace_tree (ptr2, TREE_VEC_ELT (it, 0),
+ TREE_VEC_ELT (it, 1));
+ }
OMP_CLAUSE_SIZE (node3) = fold_build2 (MINUS_EXPR, ptrdiff_type_node,
ptr, ptr2);
}
@@ -4306,6 +4315,9 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
}
break;
case OMP_LIST_MAP:
+ iterator = NULL_TREE;
+ prev = NULL;
+ prev_clauses = omp_clauses;
for (; n != NULL; n = n->next)
{
if (!openacc)
@@ -4341,6 +4353,33 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
|| n->sym->attr.flavor == FL_PARAMETER)
continue;
+ if (iterator && prev->u2.ns != n->u2.ns)
+ {
+ /* Finish previous iterator group. */
+ BLOCK_SUBBLOCKS (tree_block) = gfc_finish_block (&iter_block);
+ TREE_VEC_ELT (iterator, 5) = tree_block;
+ for (tree c = omp_clauses; c != prev_clauses;
+ c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && OMP_CLAUSE_MAP_KIND (c) !=
GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ OMP_CLAUSE_ITERATORS (c) = iterator;
+ prev_clauses = omp_clauses;
+ iterator = NULL_TREE;
+ }
+ if (n->u2.ns && (!prev || prev->u2.ns != n->u2.ns))
+ {
+ /* Start a new iterator group. */
+ gfc_init_block (&iter_block);
+ tree_block = make_node (BLOCK);
+ TREE_USED (tree_block) = 1;
+ BLOCK_VARS (tree_block) = NULL_TREE;
+ prev_clauses = omp_clauses;
+ iterator = handle_iterator (n->u2.ns, block, tree_block);
+ }
+ if (!iterator)
+ gfc_init_block (&iter_block);
+ prev = n;
+
location_t map_loc = gfc_get_location (&n->where);
bool always_modifier = false;
tree node = build_omp_clause (map_loc, OMP_CLAUSE_MAP);
@@ -4551,7 +4590,7 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
TRUTH_NOT_EXPR,
boolean_type_node,
present);
- gfc_add_expr_to_block (block,
+ gfc_add_expr_to_block (&iter_block,
build3_loc (input_location,
COND_EXPR,
void_type_node,
@@ -4621,7 +4660,8 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
tree type = TREE_TYPE (decl);
tree ptr = gfc_conv_descriptor_data_get (decl);
if (present)
- ptr = gfc_build_cond_assign_expr (block, present, ptr,
+ ptr = gfc_build_cond_assign_expr (&iter_block,
+ present, ptr,
null_pointer_node);
gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
ptr = build_fold_indirect_ref (ptr);
@@ -4648,7 +4688,7 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
ptr = gfc_conv_descriptor_data_get (decl);
ptr = gfc_build_addr_expr (NULL, ptr);
ptr = gfc_build_cond_assign_expr (
- block, present, ptr, null_pointer_node);
+ &iter_block, present, ptr, null_pointer_node);
ptr = build_fold_indirect_ref (ptr);
OMP_CLAUSE_DECL (node3) = ptr;
}
@@ -4737,7 +4777,7 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
TRUTH_ANDIF_EXPR,
boolean_type_node,
present, cond);
- gfc_add_expr_to_block (block,
+ gfc_add_expr_to_block (&iter_block,
build3_loc (input_location,
COND_EXPR,
void_type_node,
@@ -4766,12 +4806,12 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
tree cond = build3_loc (input_location, COND_EXPR,
void_type_node, present,
cond_body, NULL_TREE);
- gfc_add_expr_to_block (block, cond);
+ gfc_add_expr_to_block (&iter_block, cond);
OMP_CLAUSE_SIZE (node) = var;
}
else
{
- gfc_add_block_to_block (block, &cond_block);
+ gfc_add_block_to_block (&iter_block, &cond_block);
OMP_CLAUSE_SIZE (node) = size;
}
}
@@ -4783,8 +4823,9 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
/* A single indirectref is handled by the middle end. */
gcc_assert (!POINTER_TYPE_P (TREE_TYPE (decl)));
tree tmp = TREE_OPERAND (decl, 0);
- tmp = gfc_build_cond_assign_expr (block, present, tmp,
- null_pointer_node);
+ tmp = gfc_build_cond_assign_expr (&iter_block,
+ present, tmp,
+ null_pointer_node);
OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (tmp);
}
else
@@ -4817,7 +4858,7 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
size_type_node,
cond, size,
size_zero_node);
- size = gfc_evaluate_now (size, block);
+ size = gfc_evaluate_now (size, &iter_block);
OMP_CLAUSE_SIZE (node) = size;
}
if ((TREE_CODE (decl) != PARM_DECL
@@ -4833,7 +4874,8 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
tmp = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
: TYPE_SIZE_UNIT (TREE_TYPE (decl));
tree var = gfc_create_var (TREE_TYPE (tmp), NULL);
- gfc_add_modify_loc (input_location, block, var, tmp);
+ gfc_add_modify_loc (input_location, &iter_block,
+ var, tmp);
OMP_CLAUSE_SIZE (node) = var;
gfc_allocate_lang_decl (var);
if (TREE_CODE (decl) == INDIRECT_REF)
@@ -4863,9 +4905,10 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
&& !(POINTER_TYPE_P (type)
&& GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (type))))
k = GOMP_MAP_FIRSTPRIVATE_POINTER;
- gfc_trans_omp_array_section (block, op, n, decl, element,
- !openacc, k, node, node2,
- node3, node4);
+ gfc_trans_omp_array_section (&iter_block, op, n, decl,
+ element, !openacc, k,
+ node, node2, node3, node4,
+ iterator);
}
else if (n->expr
&& n->expr->expr_type == EXPR_VARIABLE
@@ -4881,12 +4924,12 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
gfc_init_se (&se, NULL);
gfc_conv_expr (&se, n->expr);
- gfc_add_block_to_block (block, &se.pre);
+ gfc_add_block_to_block (&iter_block, &se.pre);
/* For BT_CHARACTER a pointer is returned. */
OMP_CLAUSE_DECL (node)
= POINTER_TYPE_P (TREE_TYPE (se.expr))
? build_fold_indirect_ref (se.expr) : se.expr;
- gfc_add_block_to_block (block, &se.post);
+ gfc_add_block_to_block (&iter_block, &se.post);
if (pointer || allocatable)
{
/* If it's a bare attach/detach clause, we just want
@@ -4947,7 +4990,8 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
? DECL_SIZE_UNIT (se.expr)
: TYPE_SIZE_UNIT (TREE_TYPE (se.expr)));
tree var = gfc_create_var (TREE_TYPE (tmp), NULL);
- gfc_add_modify_loc (input_location, block, var, tmp);
+ gfc_add_modify_loc (input_location, &iter_block,
+ var, tmp);
OMP_CLAUSE_SIZE (node) = var;
gfc_allocate_lang_decl (var);
if (TREE_CODE (se.expr) == INDIRECT_REF)
@@ -5106,7 +5150,8 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
to ensure that it is not gimplified + is a decl.
*/
tree tmp = OMP_CLAUSE_SIZE (node);
tree var = gfc_create_var (TREE_TYPE (tmp), NULL);
- gfc_add_modify_loc (input_location, block, var, tmp);
+ gfc_add_modify_loc (input_location, &iter_block,
+ var, tmp);
OMP_CLAUSE_SIZE (node) = var;
gfc_allocate_lang_decl (var);
if (TREE_CODE (inner) == INDIRECT_REF)
@@ -5139,7 +5184,8 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
{
bool drop_mapping = gfc_map_array_descriptor (
node, node2, node3, node4, inner, openacc, map_loc,
- block, op, n, sym_rooted_nl, se, clauses, false);
+ &iter_block, op, n, sym_rooted_nl, se, clauses,
+ false);
if (drop_mapping)
continue;
}
@@ -5151,9 +5197,10 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
/* An array element or section. */
bool element = lastref->u.ar.type == AR_ELEMENT;
gomp_map_kind kind = GOMP_MAP_ATTACH_DETACH;
- gfc_trans_omp_array_section (block, op, n, inner, element,
- !openacc, kind, node, node2,
- node3, node4);
+ gfc_trans_omp_array_section (&iter_block, op, n, inner,
+ element, !openacc, kind,
+ node, node2, node3, node4,
+ iterator);
}
else
gcc_unreachable ();
@@ -5171,7 +5218,8 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
tree node4 = NULL_TREE;
gfc_map_array_descriptor (node1, node2, node3, node4,
mid_descr[i], openacc,
- map_loc, block, op, n,
+ map_loc, &iter_block,
+ op, n,
sym_rooted_nl, se, clauses,
true);
@@ -5194,6 +5242,9 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
finalize_map_clause:
+ if (!iterator)
+ gfc_add_block_to_block (block, &iter_block);
+
omp_clauses = gfc_trans_add_clause (node, omp_clauses);
if (node2)
omp_clauses = gfc_trans_add_clause (node2, omp_clauses);
@@ -5204,6 +5255,17 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
if (node5)
omp_clauses = gfc_trans_add_clause (node5, omp_clauses);
}
+ if (iterator)
+ {
+ /* Finish last iterator group. */
+ BLOCK_SUBBLOCKS (tree_block) = gfc_finish_block (&iter_block);
+ TREE_VEC_ELT (iterator, 5) = tree_block;
+ for (tree c = omp_clauses; c != prev_clauses;
+ c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && OMP_CLAUSE_MAP_KIND (c) !=
GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ OMP_CLAUSE_ITERATORS (c) = iterator;
+ }
break;
case OMP_LIST_TO:
case OMP_LIST_FROM:
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 0bd1df640f8..b1f263a0826 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -75,6 +75,7 @@ along with GCC; see the file COPYING3. If not see
#include "recog.h"
#include "output.h"
#include "gimplify_reg_info.h"
+#include "tree-ssa-loop-niter.h" /* For simplify_replace_tree. */
/* Identifier for a basic condition, mapping it to other basic conditions of
its Boolean expression. Basic conditions given the same uid (in the same
@@ -9744,10 +9745,17 @@ compute_omp_iterator_count (tree it, gimple_seq *pre_p)
endmbegin = fold_build2_loc (loc, POINTER_DIFF_EXPR, stype, end, begin);
else
endmbegin = fold_build2_loc (loc, MINUS_EXPR, type, end, begin);
- tree stepm1 = fold_build2_loc (loc, MINUS_EXPR, stype, step,
- build_int_cst (stype, 1));
- tree stepp1 = fold_build2_loc (loc, PLUS_EXPR, stype, step,
- build_int_cst (stype, 1));
+ /* Account for iteration stopping on the end value in Fortran rather
+ than before it. */
+ tree stepm1 = step;
+ tree stepp1 = step;
+ if (!lang_GNU_Fortran ())
+ {
+ stepm1 = fold_build2_loc (loc, MINUS_EXPR, stype, step,
+ build_int_cst (stype, 1));
+ stepp1 = fold_build2_loc (loc, PLUS_EXPR, stype, step,
+ build_int_cst (stype, 1));
+ }
tree pos = fold_build2_loc (loc, PLUS_EXPR, stype,
unshare_expr (endmbegin), stepm1);
pos = fold_build2_loc (loc, TRUNC_DIV_EXPR, stype, pos, step);
@@ -9759,10 +9767,11 @@ compute_omp_iterator_count (tree it, gimple_seq *pre_p)
}
neg = fold_build2_loc (loc, TRUNC_DIV_EXPR, stype, neg, step);
step = NULL_TREE;
- tree cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node, begin,
end);
+ tree_code cmp_op = lang_GNU_Fortran () ? LE_EXPR : LT_EXPR;
+ tree cond = fold_build2_loc (loc, cmp_op, boolean_type_node, begin, end);
pos = fold_build3_loc (loc, COND_EXPR, stype, cond, pos,
build_int_cst (stype, 0));
- cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node, end, begin);
+ cond = fold_build2_loc (loc, cmp_op, boolean_type_node, end, begin);
neg = fold_build3_loc (loc, COND_EXPR, stype, cond, neg,
build_int_cst (stype, 0));
tree osteptype = TREE_TYPE (orig_step);
@@ -9791,6 +9800,7 @@ build_omp_iterator_loop (tree it, gimple_seq *pre_p, tree
*last_bind)
if (*last_bind)
gimplify_and_add (*last_bind, pre_p);
tree block = TREE_VEC_ELT (it, 5);
+ tree block_stmts = lang_GNU_Fortran () ? BLOCK_SUBBLOCKS (block) : NULL_TREE;
*last_bind = build3 (BIND_EXPR, void_type_node,
BLOCK_VARS (block), NULL, block);
TREE_SIDE_EFFECTS (*last_bind) = 1;
@@ -9802,6 +9812,7 @@ build_omp_iterator_loop (tree it, gimple_seq *pre_p, tree
*last_bind)
tree end = TREE_VEC_ELT (it, 2);
tree step = TREE_VEC_ELT (it, 3);
tree orig_step = TREE_VEC_ELT (it, 4);
+ block = TREE_VEC_ELT (it, 5);
tree type = TREE_TYPE (var);
location_t loc = DECL_SOURCE_LOCATION (var);
/* Emit:
@@ -9812,9 +9823,9 @@ build_omp_iterator_loop (tree it, gimple_seq *pre_p, tree
*last_bind)
var = var + step;
cond_label:
if (orig_step > 0) {
- if (var < end) goto beg_label;
+ if (var < end) goto beg_label; // <= for Fortran
} else {
- if (var > end) goto beg_label;
+ if (var > end) goto beg_label; // >= for Fortran
}
for each iterator, with inner iterators added to
the ... above. */
@@ -9840,10 +9851,12 @@ build_omp_iterator_loop (tree it, gimple_seq *pre_p,
tree *last_bind)
append_to_statement_list_force (tem, p);
tem = build1 (LABEL_EXPR, void_type_node, cond_label);
append_to_statement_list (tem, p);
- tree cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node, var, end);
+ tree cond = fold_build2_loc (loc, lang_GNU_Fortran () ? LE_EXPR :
LT_EXPR,
+ boolean_type_node, var, end);
tree pos = fold_build3_loc (loc, COND_EXPR, void_type_node, cond,
build_and_jump (&beg_label), void_node);
- cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node, var, end);
+ cond = fold_build2_loc (loc, lang_GNU_Fortran () ? GE_EXPR : GT_EXPR,
+ boolean_type_node, var, end);
tree neg = fold_build3_loc (loc, COND_EXPR, void_type_node, cond,
build_and_jump (&beg_label), void_node);
tree osteptype = TREE_TYPE (orig_step);
@@ -9852,6 +9865,11 @@ build_omp_iterator_loop (tree it, gimple_seq *pre_p,
tree *last_bind)
tem = fold_build3_loc (loc, COND_EXPR, void_type_node, cond, pos, neg);
append_to_statement_list_force (tem, p);
p = &BIND_EXPR_BODY (bind);
+ /* The Fortran front-end stashes statements into the BLOCK_SUBBLOCKS
+ of the last element of the first iterator. These should go into the
+ body of the innermost loop. */
+ if (!TREE_CHAIN (it))
+ append_to_statement_list_force (block_stmts, p);
}
return p;
@@ -9985,8 +10003,14 @@ remove_unused_omp_iterator_vars (tree *list_p)
i++;
}
}
+ tree old_block = TREE_VEC_ELT (OMP_CLAUSE_ITERATORS (c), 5);
tree new_block = make_node (BLOCK);
BLOCK_VARS (new_block) = new_vars;
+ if (BLOCK_SUBBLOCKS (old_block))
+ {
+ BLOCK_SUBBLOCKS (new_block) = BLOCK_SUBBLOCKS (old_block);
+ BLOCK_SUBBLOCKS (old_block) = NULL_TREE;
+ }
TREE_VEC_ELT (new_iters, 5) = new_block;
new_iterators.safe_push (new_iters);
iter_vars.safe_push (vars.copy ());
@@ -10662,7 +10686,8 @@ build_omp_struct_comp_nodes (enum tree_code code, tree
grp_start, tree grp_end,
static tree
extract_base_bit_offset (tree base, poly_int64 *bitposp,
poly_offset_int *poffsetp,
- bool *variable_offset)
+ bool *variable_offset,
+ tree iterator)
{
tree offset;
poly_int64 bitsize, bitpos;
@@ -10672,6 +10697,18 @@ extract_base_bit_offset (tree base, poly_int64
*bitposp,
STRIP_NOPS (base);
+ if (iterator)
+ {
+ /* Replace any iterator variables with the lower bound of the iterator.
+ This will give us the nominal offset and bit position of the first
+ element, which is all we should need to lay out the mappings. The
+ actual locations of the iterated mappings are elsewhere. */
+ tree it;
+ for (it = iterator; it; it = TREE_CHAIN (it))
+ base = simplify_replace_tree (base, TREE_VEC_ELT (it, 0),
+ TREE_VEC_ELT (it, 1));
+ }
+
base = get_inner_reference (base, &bitsize, &bitpos, &offset, &mode,
&unsignedp, &reversep, &volatilep);
@@ -12498,8 +12535,11 @@ omp_accumulate_sibling_list (enum omp_region_type
region_type,
}
bool variable_offset;
+ tree iterators = OMP_CLAUSE_HAS_ITERATORS (grp_end)
+ ? OMP_CLAUSE_ITERATORS (grp_end) : NULL_TREE;
tree base
- = extract_base_bit_offset (ocd, &cbitpos, &coffset, &variable_offset);
+ = extract_base_bit_offset (ocd, &cbitpos, &coffset, &variable_offset,
+ iterators);
int base_token;
for (base_token = addr_tokens.length () - 1; base_token >= 0; base_token--)
@@ -12828,8 +12868,12 @@ omp_accumulate_sibling_list (enum omp_region_type
region_type,
sc_decl = TREE_OPERAND (sc_decl, 0);
bool variable_offset2;
+ tree iterators2 = OMP_CLAUSE_HAS_ITERATORS (*sc)
+ ? OMP_CLAUSE_ITERATORS (*sc) : NULL_TREE;
+
tree base2 = extract_base_bit_offset (sc_decl, &bitpos, &offset,
- &variable_offset2);
+ &variable_offset2,
+ iterators2);
if (!base2 || !operand_equal_p (base2, base, 0))
break;
if (scp)
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-1.f90
b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-1.f90
new file mode 100644
index 00000000000..25abbaf741e
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-1.f90
@@ -0,0 +1,26 @@
+! { dg-do compile }
+! { dg-options "-fopenmp" }
+
+program main
+ implicit none
+
+ integer, parameter :: DIM1 = 17
+ integer, parameter :: DIM2 = 39
+ type :: array_ptr
+ integer, pointer :: ptr(:)
+ end type
+
+ type (array_ptr) :: x(DIM1), y(DIM1)
+
+ !$omp target map (iterator(i=1:DIM1), to: x(i)%ptr(:))
+ !$omp end target
+
+ !$omp target map (iterator(i=1:DIM1), to: x(i)%ptr(:), y(i)%ptr(:))
+ !$omp end target
+
+ !$omp target map (iterator(i=1:DIM1), to: x(i)%ptr(:) + 3) ! { dg-error
"Syntax error in OpenMP variable list at .1." }
+ !$omp end target ! { dg-error "Unexpected \\\!\\\$OMP END TARGET statement
at .1." }
+
+ !$omp target map(iterator(i=1:DIM1), iterator(j=1:DIM2), to: x(i)%ptr(j)) !
{ dg-error "too many 'iterator' modifiers at .1." }
+ !$omp end target ! { dg-error "Unexpected \\\!\\\$OMP END TARGET statement
at .1." }
+end program
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-2.f90
b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-2.f90
new file mode 100644
index 00000000000..59c1d739cbf
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-2.f90
@@ -0,0 +1,30 @@
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-gimple" }
+
+program main
+ implicit none
+
+ integer, parameter :: DIM = 40
+ type :: array_ptr
+ integer, pointer :: ptr(:)
+ end type
+
+ type (array_ptr) :: x(DIM), y(DIM), z(DIM)
+
+ !$omp target map(iterator(i=1:10), to: x)
+ ! Add a reference to x to ensure that the 'to' clause does not get dropped.
+ x(1)%ptr(1) = 0
+ !$omp end target
+
+ !$omp target map(iterator(i2=1:10, j2=1:20), from: x(i2))
+ !$omp end target
+
+ !$omp target map(iterator(i3=1:10, j3=1:20, k3=1:30), to: x(i3+j3),
y(j3+k3), z(k3+i3))
+ !$omp end target
+end program
+
+! { dg-final { scan-tree-dump-times "map\\\(to:x" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\\(iterator\\\(integer\\\(kind=4\\\)
i2=1:10:1, loop_label=\[^\\\)\]+\\\):from:" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\\(iterator\\\(integer\\\(kind=4\\\)
j3=1:20:1, integer\\\(kind=4\\\) i3=1:10:1, loop_label=\[^\\\)\]+\\\):to:" 1
"gimple" } }
+! { dg-final { scan-tree-dump-times "map\\\(iterator\\\(integer\\\(kind=4\\\)
k3=1:30:1, integer\\\(kind=4\\\) j3=1:20:1, loop_label=\[^\\\)\]+\\\):to:" 1
"gimple" } }
+! { dg-final { scan-tree-dump-times "map\\\(iterator\\\(integer\\\(kind=4\\\)
k3=1:30:1, integer\\\(kind=4\\\) i3=1:10:1, loop_label=\[^\\\)\]+\\\):to:" 1
"gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90
b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90
new file mode 100644
index 00000000000..7dad2a69c61
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90
@@ -0,0 +1,24 @@
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-gimple" }
+
+program main
+ implicit none
+
+ integer, parameter :: DIM1 = 17
+ integer, parameter :: DIM2 = 27
+ type :: ptr_t
+ integer, pointer :: ptr(:)
+ end type
+
+ type (ptr_t) :: x(DIM1), y(DIM2)
+
+ !$omp target map(iterator(i=1:DIM1), to: x(i)%ptr(:))
map(iterator(i=1:DIM2), from: y(i)%ptr(:))
+ !$omp end target
+end program
+
+! { dg-final { scan-tree-dump-times "if \\(i <= 17\\) goto <D\\\.\[0-9\]+>;
else goto <D\\\.\[0-9\]+>;" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "if \\(i <= 27\\) goto <D\\\.\[0-9\]+>;
else goto <D\\\.\[0-9\]+>;" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\)
i=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+,
index=D\\.\[0-9\]+\\):to:MEM <\[^>\]+> \\\[\\\(\[^ \]+
\\\*\\\)D\\\.\[0-9\]+\\\]" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\)
i=1:27:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+,
index=D\\.\[0-9\]+\\):from:MEM <\[^>\]+> \\\[\\\(\[^ \]+
\\\*\\\)D\\\.\[0-9\]+\\\]" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\)
i=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+,
index=D\\.\[0-9\]+\\):attach:x\\\[D\\\.\[0-9\]+\\\]\.ptr\.data" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\)
i=1:27:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+,
index=D\\.\[0-9\]+\\):attach:y\\\[D\\\.\[0-9\]+\\\]\.ptr\.data" 1 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-4.f90
b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-4.f90
new file mode 100644
index 00000000000..3359eaad0b4
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-4.f90
@@ -0,0 +1,28 @@
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-gimple" }
+
+module m
+ !$omp declare target (baz)
+ interface
+ subroutine baz (x, p)
+ integer, intent(in) :: x
+ integer, pointer :: p(:)
+ end subroutine
+ integer function bar (x, i)
+ integer :: x, i
+ end function
+ end interface
+contains
+ subroutine foo (x, p)
+ integer :: x
+ integer, pointer :: p(:)
+
+ !$omp target map (iterator (i=1:4), to: p(bar (x, i)))
+ call baz (x, p)
+ !$omp end target
+ end subroutine
+end module
+
+! { dg-final { scan-tree-dump "firstprivate\\\(x\\\)" "gimple" } }
+! { dg-final { scan-tree-dump-times "bar \\\(x, &" 2 "gimple" } }
+! { dg-final { scan-tree-dump "map\\\(iterator\\\(integer\\\(kind=4\\\)
i=1:4:1, loop_label=" "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-6.f90
b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-6.f90
new file mode 100644
index 00000000000..b83481920b6
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-6.f90
@@ -0,0 +1,26 @@
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-omplower" }
+
+module m
+ implicit none
+ integer, parameter :: DIM = 17
+ type :: array_ptr
+ integer, pointer :: ptr(:)
+ end type
+contains
+ subroutine f(x)
+ type (array_ptr) :: x(DIM)
+ integer :: i
+
+ !$omp target data map(iterator(it = 1:DIM), tofrom: x(it)%ptr(:))
+ do i = 1, DIM
+ x(i)%ptr(1) = i
+ end do
+ !$omp end target data
+ end subroutine
+end module
+
+! { dg-final { scan-tree-dump "if \\(it <= 17\\) goto <D\\\.\[0-9\]+>; else
goto <D\\\.\[0-9\]+>;" "omplower" } }
+! { dg-final { scan-tree-dump "map\\(iterator\\(integer\\(kind=\[0-9\]+\\)
it=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+,
index=D\\\.\[0-9\]+\\):to:\\(\\\*x\\\.\[0-9\]+\\)\\\[D\\\.\[0-9\]+\\\]\\\.ptr"
"omplower" } }
+! { dg-final { scan-tree-dump "map\\(iterator\\(integer\\(kind=\[0-9\]+\\)
it=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+,
index=D\\\.\[0-9\]+\\):tofrom:MEM" "omplower" } }
+! { dg-final { scan-tree-dump "map\\(iterator\\(integer\\(kind=\[0-9\]+\\)
it=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+,
index=D\\\.\[0-9\]+\\):attach:\\(\\\*x\\\.\[0-9\]+\\)\\\[D\\\.\[0-9\]+\\\]\\\.ptr.data"
"omplower" } }
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 79e0a734dbb..0674a27501b 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1848,7 +1848,9 @@ dump_block_node (pretty_printer *pp, tree block, int spc,
dump_flags_t flags)
newline_and_indent (pp, spc + 2);
}
- if (BLOCK_SUBBLOCKS (block))
+ if (BLOCK_SUBBLOCKS (block)
+ && (!lang_GNU_Fortran ()
+ || TREE_CODE (BLOCK_SUBBLOCKS (block)) != STATEMENT_LIST))
{
pp_string (pp, "SUBBLOCKS: ");
for (t = BLOCK_SUBBLOCKS (block); t; t = BLOCK_CHAIN (t))
diff --git a/libgomp/target.c b/libgomp/target.c
index 7543065d106..6122f2cd2a0 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1060,6 +1060,41 @@ kind_to_name (unsigned short kind, bool short_mapkind)
}
}
+static void
+gomp_add_map (size_t idx, size_t *new_idx,
+ void ***hostaddrs, size_t **sizes, unsigned short **skinds,
+ void ***new_hostaddrs, size_t **new_sizes,
+ unsigned short **new_kinds, size_t *iterator_count)
+{
+ if ((*sizes)[idx] == SIZE_MAX)
+ {
+ uintptr_t *iterator_array = (*hostaddrs)[idx];
+ size_t count = *iterator_array++;
+ for (size_t i = 0; i < count; i++)
+ {
+ (*new_hostaddrs)[*new_idx] = (void *) *iterator_array++;
+ (*new_sizes)[*new_idx] = *iterator_array++;
+ (*new_kinds)[*new_idx] = (*skinds)[idx];
+ iterator_count[*new_idx] = i + 1;
+ gomp_debug (1,
+ "Expanding map %u <%s>: "
+ "hostaddrs[%u] = %p, sizes[%u] = %lu\n",
+ (int) idx, kind_to_name ((*new_kinds)[*new_idx], true),
+ (int) *new_idx, (*new_hostaddrs)[*new_idx],
+ (int) *new_idx, (unsigned long) (*new_sizes)[*new_idx]);
+ (*new_idx)++;
+ }
+ }
+ else
+ {
+ (*new_hostaddrs)[*new_idx] = (*hostaddrs)[idx];
+ (*new_sizes)[*new_idx] = (*sizes)[idx];
+ (*new_kinds)[*new_idx] = (*skinds)[idx];
+ iterator_count[*new_idx] = 0;
+ (*new_idx)++;
+ }
+}
+
/* When GCC encounters a clause with an iterator, e.g.:
#pragma omp target map (iterator(i=0:4), to: x[i])
@@ -1093,6 +1128,7 @@ gomp_merge_iterator_maps (size_t *mapnum, void
***hostaddrs, size_t **sizes,
bool iterator_p = false;
size_t map_count = 0;
unsigned short **skinds = (unsigned short **) kinds;
+ const int typemask = 0xff;
for (size_t i = 0; i < *mapnum; i++)
if ((*sizes)[i] == SIZE_MAX)
@@ -1119,33 +1155,35 @@ gomp_merge_iterator_maps (size_t *mapnum, void
***hostaddrs, size_t **sizes,
for (size_t i = 0; i < *mapnum; i++)
{
- if ((*sizes)[i] == SIZE_MAX)
+ int map_type = get_kind (true, *skinds, i) & typemask;
+ if (map_type == GOMP_MAP_STRUCT || map_type == GOMP_MAP_STRUCT_UNORD)
{
- uintptr_t *iterator_array = (*hostaddrs)[i];
- size_t count = *iterator_array++;
- for (size_t j = 0; j < count; j++)
+ size_t field_count = (*sizes)[i];
+ size_t idx_i = new_idx;
+
+ gomp_add_map (i, &new_idx, hostaddrs, sizes, skinds,
+ &new_hostaddrs, &new_sizes, &new_kinds,
+ *iterator_count);
+
+ for (size_t j = i + 1; j <= i + field_count; j++)
{
- new_hostaddrs[new_idx] = (void *) *iterator_array++;
- new_sizes[new_idx] = *iterator_array++;
- new_kinds[new_idx] = (*skinds)[i];
- (*iterator_count)[new_idx] = j + 1;
- gomp_debug (1,
- "Expanding map %u <%s>: "
- "hostaddrs[%u] = %p, sizes[%u] = %lu\n",
- (int) i, kind_to_name (new_kinds[new_idx], true),
- (int) new_idx, new_hostaddrs[new_idx],
- (int) new_idx, (unsigned long) new_sizes[new_idx]);
- new_idx++;
+ if ((*sizes)[j] == SIZE_MAX)
+ {
+ uintptr_t *iterator_array = (*hostaddrs)[j];
+ size_t count = iterator_array[0];
+ new_sizes[idx_i] += count - 1;
+ }
+ gomp_add_map (j, &new_idx, hostaddrs, sizes, skinds,
+ &new_hostaddrs, &new_sizes, &new_kinds,
+ *iterator_count);
}
+ gomp_debug (1, "Map %u: new field count = %lu\n",
+ (int) i, (unsigned long) new_sizes[idx_i]);
+ i += field_count;
}
else
- {
- new_hostaddrs[new_idx] = (*hostaddrs)[i];
- new_sizes[new_idx] = (*sizes)[i];
- new_kinds[new_idx] = (*skinds)[i];
- (*iterator_count)[new_idx] = 0;
- new_idx++;
- }
+ gomp_add_map (i, &new_idx, hostaddrs, sizes, skinds,
+ &new_hostaddrs, &new_sizes, &new_kinds, *iterator_count);
}
*mapnum = map_count;
diff --git a/libgomp/testsuite/libgomp.fortran/target-map-iterators-1.f90
b/libgomp/testsuite/libgomp.fortran/target-map-iterators-1.f90
new file mode 100644
index 00000000000..80e077e69fd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-map-iterators-1.f90
@@ -0,0 +1,45 @@
+! { dg-do run }
+
+! Test transfer of dynamically-allocated arrays to target using map
+! iterators.
+
+program test
+ implicit none
+
+ integer, parameter :: DIM1 = 8
+ integer, parameter :: DIM2 = 15
+
+ type :: array_ptr
+ integer, pointer :: arr(:)
+ end type
+
+ type (array_ptr) :: x(DIM1)
+ integer :: expected, sum, i, j
+
+ expected = mkarray ()
+
+ !$omp target map(iterator(i=1:DIM1), to: x(i)%arr(:)) map(from: sum)
+ sum = 0
+ do i = 1, DIM1
+ do j = 1, DIM2
+ sum = sum + x(i)%arr(j)
+ end do
+ end do
+ !$omp end target
+
+ if (sum .ne. expected) stop 1
+contains
+ integer function mkarray ()
+ integer :: exp = 0
+
+ do i = 1, DIM1
+ allocate (x(i)%arr(DIM2))
+ do j = 1, DIM2
+ x(i)%arr(j) = i * j
+ exp = exp + x(i)%arr(j)
+ end do
+ end do
+
+ mkarray = exp
+ end function
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/target-map-iterators-2.f90
b/libgomp/testsuite/libgomp.fortran/target-map-iterators-2.f90
new file mode 100644
index 00000000000..cf0e7fbd9b3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-map-iterators-2.f90
@@ -0,0 +1,45 @@
+! { dg-do run }
+
+! Test transfer of dynamically-allocated arrays from target using map
+! iterators.
+
+program test
+ implicit none
+
+ integer, parameter :: DIM1 = 8
+ integer, parameter :: DIM2 = 15
+
+ type :: array_ptr
+ integer, pointer :: arr(:)
+ end type
+
+ type (array_ptr) :: x(DIM1)
+ integer :: expected, sum, i, j
+
+ call mkarray
+
+ !$omp target map(iterator(i=1:DIM1), from: x(i)%arr(:)) map(from: expected)
+ expected = 0
+ do i = 1, DIM1
+ do j = 1, DIM2
+ x(i)%arr(j) = (i+1) * (j+1)
+ expected = expected + x(i)%arr(j)
+ end do
+ end do
+ !$omp end target
+
+ sum = 0
+ do i = 1, DIM1
+ do j = 1, DIM2
+ sum = sum + x(i)%arr(j)
+ end do
+ end do
+
+ if (sum .ne. expected) stop 1
+contains
+ subroutine mkarray
+ do i = 1, DIM1
+ allocate (x(i)%arr(DIM2))
+ end do
+ end subroutine
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/target-map-iterators-3.f90
b/libgomp/testsuite/libgomp.fortran/target-map-iterators-3.f90
new file mode 100644
index 00000000000..b7ba912a886
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-map-iterators-3.f90
@@ -0,0 +1,50 @@
+! { dg-do run }
+
+! Test transfer of dynamically-allocated arrays to target using map
+! iterators, with multiple iterators and function calls in the iterator
+! expression.
+
+program test
+ implicit none
+
+ integer, parameter :: DIM1 = 16
+ integer, parameter :: DIM2 = 4
+
+ type :: array_ptr
+ integer, pointer :: arr(:)
+ end type
+
+ type (array_ptr) :: x(DIM1), y(DIM1)
+ integer :: expected, sum, i, j, k
+
+ expected = mkarrays ()
+
+ !$omp target map(iterator(i=0:DIM1/4-1, j=0:3), to: x(i * 4 + j + 1)%arr(:))
&
+ !$omp map(iterator(k=1:DIM1), to: y(k)%arr(:)) &
+ !$omp map(from: sum)
+ sum = 0
+ do i = 1, DIM1
+ do j = 1, DIM2
+ sum = sum + x(i)%arr(j) * y(i)%arr(j)
+ end do
+ end do
+ !$omp end target
+
+ if (sum .ne. expected) stop 1
+contains
+ integer function mkarrays ()
+ integer :: exp = 0
+
+ do i = 1, DIM1
+ allocate (x(i)%arr(DIM2))
+ allocate (y(i)%arr(DIM2))
+ do j = 1, DIM2
+ x(i)%arr(j) = i * j
+ y(i)%arr(j) = i + j
+ exp = exp + x(i)%arr(j) * y(i)%arr(j)
+ end do
+ end do
+
+ mkarrays = exp
+ end function
+end program
--
2.39.5
