From: Kwok Cheung Yeung <[email protected]>
This patch enables support for using non-constant expressions when specifying
iterators in the map clause of target constructs and to/from clauses of
target update constructs.
gcc/
* gimplify.cc (omp_iterator_elems_length): New.
(build_omp_iterators_loops): Change type of elements
array to pointer of pointers if array length is non-constant, and
assign size with indirect reference. Reorder elements added to
iterator vector and add element containing the iteration count. Use
omp_iterator_elems_length to compute element array size required.
* gimplify.h (omp_iterator_elems_length): New prototype.
* omp-low.cc (lower_omp_map_iterator_expr): Reorder elements read
from iterator vector. If elements field is a pointer type, assign
using pointer arithmetic followed by indirect reference, and return
the field directly.
(lower_omp_map_iterator_size): Reorder elements read from iterator
vector. If elements field is a pointer type, assign using pointer
arithmetic followed by indirect reference.
(allocate_omp_iterator_elems): New.
(free_omp_iterator_elems): New.
(lower_omp_target): Call allocate_omp_iterator_elems before inserting
loops sequence, and call free_omp_iterator_elems afterwards.
* tree-pretty-print.cc (dump_omp_iterators): Print extra elements in
iterator vector.
gcc/testsuite/
* c-c++-common/gomp/target-map-iterators-3.c: Update expected Gimple
output.
* c-c++-common/gomp/target-map-iterators-5.c: New.
* c-c++-common/gomp/target-update-iterators-3.c: Update expected
Gimple output.
* gfortran.dg/gomp/target-map-iterators-3.f90: Likewise.
* gfortran.dg/gomp/target-map-iterators-5.f90: New.
* gfortran.dg/gomp/target-map-iterators-6.f90: Update expected
Gimple output.
* gfortran.dg/gomp/target-update-iterators-3.f90: Likewise.
libgomp/
* testsuite/libgomp.c-c++-common/target-map-iterators-1c.c: New.
* testsuite/libgomp.c-c++-common/target-map-iterators-1d.c: New.
* testsuite/libgomp.c-c++-common/target-map-iterators-1e.c: New.
* testsuite/libgomp.c-c++-common/target-map-iterators-1f.c: New.
* testsuite/libgomp.c-c++-common/target-map-iterators-4.c: New.
* testsuite/libgomp.c-c++-common/target-map-iterators-5.c: New.
* testsuite/libgomp.c-c++-common/target-update-iterators-4.c: New.
* testsuite/libgomp.fortran/target-map-iterators-1c.f90: New.
* testsuite/libgomp.fortran/target-map-iterators-4.f90: New.
* testsuite/libgomp.fortran/target-map-iterators-5.f90: New.
* testsuite/libgomp.fortran/target-update-iterators-4.f90: New.
Co-authored-by: Sandra Loosemore <[email protected]>
---
gcc/gimplify.cc | 48 +++++----
gcc/gimplify.h | 1 +
gcc/omp-low.cc | 100 ++++++++++++++++--
.../gomp/target-map-iterators-3.c | 8 +-
.../gomp/target-map-iterators-5.c | 14 +++
.../gomp/target-update-iterators-3.c | 4 +-
.../gomp/target-map-iterators-3.f90 | 8 +-
.../gomp/target-map-iterators-5.f90 | 21 ++++
.../gomp/target-map-iterators-6.f90 | 6 +-
.../gomp/target-update-iterators-3.f90 | 6 +-
gcc/tree-pretty-print.cc | 6 +-
.../target-map-iterators-1c.c | 58 ++++++++++
.../target-map-iterators-1d.c | 69 ++++++++++++
.../target-map-iterators-1e.c | 66 ++++++++++++
.../target-map-iterators-1f.c | 71 +++++++++++++
.../target-map-iterators-4.c | 48 +++++++++
.../target-map-iterators-5.c | 59 +++++++++++
.../target-update-iterators-4.c | 66 ++++++++++++
.../target-map-iterators-1c.f90 | 58 ++++++++++
.../target-map-iterators-4.f90 | 48 +++++++++
.../target-map-iterators-5.f90 | 55 ++++++++++
.../target-update-iterators-4.f90 | 70 ++++++++++++
22 files changed, 841 insertions(+), 49 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-map-iterators-5.c
create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-map-iterators-5.f90
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1c.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1d.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1e.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1f.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-4.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-5.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-4.c
create mode 100644
libgomp/testsuite/libgomp.fortran/target-map-iterators-1c.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/target-map-iterators-4.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/target-map-iterators-5.f90
create mode 100644
libgomp/testsuite/libgomp.fortran/target-update-iterators-4.f90
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index b1f263a0826..5efe59bc3ca 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -10043,6 +10043,17 @@ struct iterator_loop_info_t
typedef hash_map<tree, iterator_loop_info_t> iterator_loop_info_map_t;
+/* Returns the length of the OpenMP iterator elements array
+ for an iterator with COUNT iterations. There is one element for the
+ iteration count, and two elements (expr + size) for each iteration. */
+
+tree
+omp_iterator_elems_length (tree count)
+{
+ tree count_2 = size_binop (MULT_EXPR, count, size_int (2));
+ return size_binop (PLUS_EXPR, count_2, size_int (1));
+}
+
/* Builds a loop to expand any OpenMP iterators in the clauses in LIST_P,
reusing any previously built loops if they use the same set of iterators.
Generated Gimple statements are placed into LOOPS_SEQ_P. The clause
@@ -10101,27 +10112,21 @@ build_omp_iterators_loops (tree *list_p, gimple_seq
*loops_seq_p)
}
/* Create array to hold expanded values. */
- tree last_count_2 = size_binop (MULT_EXPR, loop.count, size_int (2));
- tree arr_length = size_binop (PLUS_EXPR, last_count_2, size_int (1));
- tree elems = NULL_TREE;
- if (TREE_CONSTANT (arr_length))
- {
- tree type = build_array_type (ptr_type_node,
- build_index_type (arr_length));
- elems = create_tmp_var_raw (type, "omp_iter_data");
- TREE_ADDRESSABLE (elems) = 1;
- gimple_add_tmp_var (elems);
- }
- else
- {
- /* Handle dynamic sizes. */
- sorry ("dynamic iterator sizes not implemented yet");
- }
+ tree arr_length = omp_iterator_elems_length (loop.count);
+ tree elems_type = TREE_CONSTANT (arr_length)
+ ? build_array_type (ptr_type_node,
+ build_index_type (arr_length))
+ : build_pointer_type (ptr_type_node);
+ tree elems = create_tmp_var_raw (elems_type, "omp_iter_data");
+ TREE_ADDRESSABLE (elems) = 1;
+ gimple_add_tmp_var (elems);
/* BEFORE LOOP: */
/* elems[0] = count; */
- tree lhs = build4 (ARRAY_REF, ptr_type_node, elems, size_int (0),
- NULL_TREE, NULL_TREE);
+ tree lhs = TREE_CODE (TREE_TYPE (elems)) == ARRAY_TYPE
+ ? build4 (ARRAY_REF, ptr_type_node, elems, size_int (0), NULL_TREE,
+ NULL_TREE)
+ : build1 (INDIRECT_REF, ptr_type_node, elems);
tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
void_type_node, lhs, loop.count);
gimplify_and_add (tem, loops_seq_p);
@@ -10129,10 +10134,11 @@ build_omp_iterators_loops (tree *list_p, gimple_seq
*loops_seq_p)
/* Make a copy of the iterator with extra info at the end. */
int elem_count = TREE_VEC_LENGTH (OMP_CLAUSE_ITERATORS (c));
tree new_iterator = copy_omp_iterator (OMP_CLAUSE_ITERATORS (c),
- elem_count + 3);
+ elem_count + 4);
TREE_VEC_ELT (new_iterator, elem_count) = loop.body_label;
- TREE_VEC_ELT (new_iterator, elem_count + 1) = elems;
- TREE_VEC_ELT (new_iterator, elem_count + 2) = loop.index;
+ TREE_VEC_ELT (new_iterator, elem_count + 1) = loop.index;
+ TREE_VEC_ELT (new_iterator, elem_count + 2) = elems;
+ TREE_VEC_ELT (new_iterator, elem_count + 3) = loop.count;
TREE_CHAIN (new_iterator) = TREE_CHAIN (OMP_CLAUSE_ITERATORS (c));
OMP_CLAUSE_ITERATORS (c) = new_iterator;
diff --git a/gcc/gimplify.h b/gcc/gimplify.h
index 18c7514c9a3..1f13f5b2dec 100644
--- a/gcc/gimplify.h
+++ b/gcc/gimplify.h
@@ -76,6 +76,7 @@ extern enum gimplify_status gimplify_expr (tree *, gimple_seq
*, gimple_seq *,
extern tree omp_get_construct_context (void);
int omp_has_novariants (void);
+extern tree omp_iterator_elems_length (tree count);
extern gimple_seq *enter_omp_iterator_loop_context (tree, gomp_target *,
gimple_seq * = NULL);
extern void exit_omp_iterator_loop_context (tree);
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index d31fa96dc8c..ce2f4b1ed93 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -12703,20 +12703,30 @@ lower_omp_map_iterator_expr (tree expr, tree c,
gomp_target *stmt)
return expr;
tree iterator = OMP_CLAUSE_ITERATORS (c);
- tree elems = TREE_VEC_ELT (iterator, 7);
- tree index = TREE_VEC_ELT (iterator, 8);
+ tree index = TREE_VEC_ELT (iterator, 7);
+ tree elems = TREE_VEC_ELT (iterator, 8);
gimple_seq *loop_body_p = enter_omp_iterator_loop_context (c, stmt);
/* IN LOOP BODY: */
/* elems[idx] = <expr>; */
- tree lhs = build4 (ARRAY_REF, ptr_type_node, elems, index,
- NULL_TREE, NULL_TREE);
+ tree lhs;
+ if (TREE_CODE (TREE_TYPE (elems)) == ARRAY_TYPE)
+ lhs = build4 (ARRAY_REF, ptr_type_node, elems, index, NULL_TREE,
NULL_TREE);
+ else
+ {
+ tree tmp = size_binop (MULT_EXPR, index, TYPE_SIZE_UNIT (ptr_type_node));
+ tmp = size_binop (POINTER_PLUS_EXPR, elems, tmp);
+ lhs = build1 (INDIRECT_REF, ptr_type_node, tmp);
+ }
tree mod_expr = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
void_type_node, lhs, expr);
gimplify_and_add (mod_expr, loop_body_p);
exit_omp_iterator_loop_context (c);
- return build_fold_addr_expr_with_type (elems, ptr_type_node);
+ if (TREE_CODE (TREE_TYPE (elems)) == ARRAY_TYPE)
+ return build_fold_addr_expr_with_type (elems, ptr_type_node);
+ else
+ return elems;
}
/* Set SIZE as the size expression that should result from the clause C
@@ -12731,15 +12741,25 @@ lower_omp_map_iterator_size (tree size, tree c,
gomp_target *stmt)
return size;
tree iterator = OMP_CLAUSE_ITERATORS (c);
- tree elems = TREE_VEC_ELT (iterator, 7);
- tree index = TREE_VEC_ELT (iterator, 8);
+ tree index = TREE_VEC_ELT (iterator, 7);
+ tree elems = TREE_VEC_ELT (iterator, 8);
gimple_seq *loop_body_p = enter_omp_iterator_loop_context (c, stmt);
/* IN LOOP BODY: */
/* elems[idx+1] = <size>; */
- tree lhs = build4 (ARRAY_REF, ptr_type_node, elems,
- size_binop (PLUS_EXPR, index, size_int (1)),
- NULL_TREE, NULL_TREE);
+ tree lhs;
+ if (TREE_CODE (TREE_TYPE (elems)) == ARRAY_TYPE)
+ lhs = build4 (ARRAY_REF, ptr_type_node, elems,
+ size_binop (PLUS_EXPR, index, size_int (1)),
+ NULL_TREE, NULL_TREE);
+ else
+ {
+ tree index_1 = size_binop (PLUS_EXPR, index, size_int (1));
+ tree tmp = size_binop (MULT_EXPR, index_1,
+ TYPE_SIZE_UNIT (ptr_type_node));
+ tmp = size_binop (POINTER_PLUS_EXPR, elems, tmp);
+ lhs = build1 (INDIRECT_REF, ptr_type_node, tmp);
+ }
tree mod_expr = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
void_type_node, lhs, size);
gimplify_and_add (mod_expr, loop_body_p);
@@ -12748,6 +12768,62 @@ lower_omp_map_iterator_size (tree size, tree c,
gomp_target *stmt)
return size_int (SIZE_MAX);
}
+static void
+allocate_omp_iterator_elems (tree clauses, gimple_seq loops_seq)
+{
+ for (tree c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (!OMP_CLAUSE_HAS_ITERATORS (c))
+ continue;
+ tree iters = OMP_CLAUSE_ITERATORS (c);
+ tree elems = TREE_VEC_ELT (iters, 8);
+ if (!POINTER_TYPE_P (TREE_TYPE (elems)))
+ continue;
+ tree arr_length = omp_iterator_elems_length (TREE_VEC_ELT (iters, 9));
+ tree call = builtin_decl_explicit (BUILT_IN_MALLOC);
+ tree size = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MULT_EXPR,
+ size_type_node, arr_length,
+ TYPE_SIZE_UNIT (ptr_type_node));
+ tree tmp = build_call_expr_loc (OMP_CLAUSE_LOCATION (c), call, 1,
+ size);
+
+ /* Find the first statement '<index> = -1' in the pre-loop statements.
*/
+ tree index = TREE_VEC_ELT (iters, 7);
+ gimple_stmt_iterator gsi;
+ for (gsi = gsi_start (loops_seq); !gsi_end_p (gsi); gsi_next (&gsi))
+ {
+ gimple *stmt = gsi_stmt (gsi);
+ if (gimple_code (stmt) == GIMPLE_ASSIGN
+ && gimple_assign_lhs (stmt) == index
+ && gimple_assign_rhs1 (stmt) == size_int (-1))
+ break;
+ }
+ gcc_assert (!gsi_end_p (gsi));
+
+ gimple_seq alloc_seq = NULL;
+ gimplify_assign (elems, tmp, &alloc_seq);
+ gsi_insert_seq_before (&gsi, alloc_seq, GSI_SAME_STMT);
+ }
+}
+
+static void
+free_omp_iterator_elems (tree clauses, gimple_seq *seq)
+{
+ for (tree c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (!OMP_CLAUSE_HAS_ITERATORS (c))
+ continue;
+ tree elems = TREE_VEC_ELT (OMP_CLAUSE_ITERATORS (c), 8);
+ if (!POINTER_TYPE_P (TREE_TYPE (elems)))
+ continue;
+ tree call = builtin_decl_explicit (BUILT_IN_FREE);
+ call = build_call_expr_loc (OMP_CLAUSE_LOCATION (c), call, 1, elems);
+ gimplify_and_add (call, seq);
+ tree clobber = build_clobber (TREE_TYPE (elems));
+ gimple_seq_add_stmt (seq, gimple_build_assign (elems, clobber));
+ }
+}
+
/* Lower the GIMPLE_OMP_TARGET in the current statement
in GSI_P. CTX holds context information for the directive. */
@@ -14602,9 +14678,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
gimple_omp_set_body (stmt, new_body);
}
+ allocate_omp_iterator_elems (clauses,
+ gimple_omp_target_iterator_loops (stmt));
gsi_insert_seq_before (gsi_p, gimple_omp_target_iterator_loops (stmt),
GSI_SAME_STMT);
gimple_omp_target_set_iterator_loops (stmt, NULL);
+ free_omp_iterator_elems (clauses, &olist);
+
bind = gimple_build_bind (NULL, NULL,
tgt_bind ? gimple_bind_block (tgt_bind)
: NULL_TREE);
diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c
b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c
index c31336538de..82927edbbf8 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c
@@ -17,7 +17,7 @@ void f (int ***x, float ***y, double **z)
/* { dg-final { scan-tree-dump-times "if \\(i <= 9\\) goto <D\\\.\[0-9\]+>;
else goto <D\\\.\[0-9\]+>;" 3 "gimple" } } */
/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto <D\\\.\[0-9\]+>;
else goto <D\\\.\[0-9\]+>;" 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1,
loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+,
index=D\\\.\[0-9\]+\\):from:\\*D\\\.\[0-9\]+" 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1,
loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+,
index=D\\\.\[0-9\]+\\):attach:\\*D\\\.\[0-9\]+" 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int
j=0:20:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+,
index=D\\\.\[0-9\]+\\):to:\\*D\\\.\[0-9\]+" 2 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int
j=0:20:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+,
index=D\\\.\[0-9\]+\\):attach:\\*D\\\.\[0-9\]+" 4 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1,
loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+,
elems=omp_iter_data\\\.\[0-9\]+, elems_count=\[0-9\]+\\):from:\\*D\\\.\[0-9\]+"
1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1,
loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+,
elems=omp_iter_data\\\.\[0-9\]+,
elems_count=\[0-9\]+\\):attach:\\*D\\\.\[0-9\]+" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int
j=0:20:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+,
elems=omp_iter_data\\\.\[0-9\]+, elems_count=\[0-9\]+\\):to:\\*D\\\.\[0-9\]+" 2
"gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int
j=0:20:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+,
elems=omp_iter_data\\\.\[0-9\]+,
elems_count=\[0-9\]+\\):attach:\\*D\\\.\[0-9\]+" 4 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-5.c
b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-5.c
new file mode 100644
index 00000000000..58daa0e73bf
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-5.c
@@ -0,0 +1,14 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-omplower" } */
+
+#define DIM2 17
+
+void f (int **x, int lbound, int ubound, int stride)
+{
+ #pragma omp target map(to:x) map(iterator(i=lbound:ubound:stride), to: x[i][
:DIM2])
+ ;
+}
+
+/* { dg-final { scan-tree-dump-times "_\[0-9\]+ = ubound - lbound;" 2
"omplower" } } */
+/* { dg-final { scan-tree-dump-times "D\\\.\[0-9\]+ = __builtin_malloc
\\(D\\\.\[0-9\]+\\);" 2 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_free
\\(omp_iter_data\\\.\[0-9\]+\\);" 2 "omplower" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c
b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c
index 0057cf76d87..00da5138e77 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c
@@ -13,5 +13,5 @@ void f (int ***x, float ***y, double **z)
/* { dg-final { scan-tree-dump-times "if \\(i <= 9\\) goto <D\.\[0-9\]+>; else
goto <D\.\[0-9\]+>;" 2 "gimple" } } */
/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto <D\.\[0-9\]+>;
else goto <D\.\[0-9\]+>;" 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "to\\(iterator\\(int i=0:10:1, int
j=0:20:1, loop_label=<D\.\[0-9\]+>, elems=omp_iter_data\.\[0-9\]+,
index=D\.\[0-9\]+\\):\\*D\.\[0-9\]+" 2 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "from\\(iterator\\(int i=0:10:1,
loop_label=<D\.\[0-9\]+>, elems=omp_iter_data\.\[0-9\]+,
index=D\.\[0-9\]+\\):\\*D\.\[0-9\]+" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "to\\(iterator\\(int i=0:10:1, int
j=0:20:1, loop_label=<D\.\[0-9\]+>, index=D\.\[0-9\]+,
elems=omp_iter_data\.\[0-9\]+, elems_count=200\\):\\*D\.\[0-9\]+" 2 "gimple" }
} */
+/* { dg-final { scan-tree-dump-times "from\\(iterator\\(int i=0:10:1,
loop_label=<D\.\[0-9\]+>, index=D\.\[0-9\]+, elems=omp_iter_data\.\[0-9\]+,
elems_count=10\\):\\*D\.\[0-9\]+" 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
index 7dad2a69c61..1099955c2cc 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90
@@ -18,7 +18,7 @@ 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" } }
+! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\)
i=1:17:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+,
elems=omp_iter_data\\\.\[0-9\]+, elems_count=17\\):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\]+>, index=D\\\.\[0-9\]+,
elems=omp_iter_data\\\.\[0-9\]+, elems_count=27\\):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\]+>, index=D\\\.\[0-9\]+,
elems=omp_iter_data\\\.\[0-9\]+,
elems_count=17\\):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\]+>, index=D\\\.\[0-9\]+,
elems=omp_iter_data\\\.\[0-9\]+,
elems_count=27\\):attach:y\\\[D\\\.\[0-9\]+\\\]\.ptr\.data" 1 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-5.f90
b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-5.f90
new file mode 100644
index 00000000000..2de312a8b9a
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-5.f90
@@ -0,0 +1,21 @@
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-omplower" }
+
+module m
+ integer, parameter :: DIM1 = 31
+ integer, parameter :: DIM2 = 17
+ type :: array_ptr
+ integer, pointer :: ptr(:)
+ end type
+contains
+ subroutine f (x, stride)
+ type (array_ptr) :: x(:)
+ integer :: stride
+
+ !$omp target map(to: x) map(iterator(i=lbound(x, 1):ubound(x, 1):stride),
to: x(i)%ptr(:))
+ !$omp end target
+ end subroutine
+end module
+
+! { dg-final { scan-tree-dump-times "D\\\.\[0-9\]+ = __builtin_malloc
\\(D\\\.\[0-9\]+\\);" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "__builtin_free
\\(omp_iter_data\\\.\[0-9\]+\\);" 2 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-6.f90
b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-6.f90
index b83481920b6..5cc36f0d024 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-6.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-6.f90
@@ -21,6 +21,6 @@ contains
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" } }
+! { dg-final { scan-tree-dump "map\\(iterator\\(integer\\(kind=\[0-9\]+\\)
it=1:17:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+,
elems=omp_iter_data\\\.\[0-9\]+,
elems_count=17\\):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\]+>, index=D\\\.\[0-9\]+,
elems=omp_iter_data\\\.\[0-9\]+, elems_count=17\\):tofrom:MEM" "omplower" } }
+! { dg-final { scan-tree-dump "map\\(iterator\\(integer\\(kind=\[0-9\]+\\)
it=1:17:1, loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+,
elems=omp_iter_data\\\.\[0-9\]+,
elems_count=17\\):attach:\\(\\\*x\\\.\[0-9\]+\\)\\\[D\\\.\[0-9\]+\\\]\\\.ptr.data"
"omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90
b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90
index d9c92cf4679..a8dffcf09ba 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90
@@ -18,6 +18,6 @@ program test
end program
! { dg-final { scan-tree-dump-times "if \\(i <= 17\\) goto <D\\\.\[0-9\]+>;
else goto <D\\\.\[0-9\]+>;" 2 "gimple" } }
-! { dg-final { scan-tree-dump-times "if \\(j <= 39\\) goto <D\\\.\[0-9\]+>;
else goto <D\\\.\[0-9\]+>;" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "to\\(iterator\\(integer\\(kind=4\\)
j=1:39:1, integer\\(kind=4\\) i=1:17:1, loop_label=<D\\\.\[0-9\]+>,
elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):MEM <\[^>\]+>
\\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]" 2 "gimple" } }
-! { dg-final { scan-tree-dump-times "from\\(iterator\\(integer\\(kind=4\\)
i=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+,
index=D\\\.\[0-9\]+\\):MEM <\[^>\]+> \\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]"
1 "gimple" } }
+! { dg-final { scan-tree-dump "if \\(j <= 39\\) goto <D\\\.\[0-9\]+>; else
goto <D\\\.\[0-9\]+>;" "gimple" } }
+! { dg-final { scan-tree-dump-times "to\\(iterator\\(integer\\(kind=4\\)
j=1:39:1, integer\\(kind=4\\) i=1:17:1, loop_label=<D\\\.\[0-9\]+>,
index=D\\\.\[0-9\]+, elems=omp_iter_data\\\.\[0-9\]+, elems_count=663\\):MEM
<\[^>\]+> \\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]" 2 "gimple" } }
+! { dg-final { scan-tree-dump "from\\(iterator\\(integer\\(kind=4\\) i=1:17:1,
loop_label=<D\\\.\[0-9\]+>, index=D\\\.\[0-9\]+,
elems=omp_iter_data\\\.\[0-9\]+, elems_count=17\\):MEM <\[^>\]+> \\\[\\\(\[^
\]+ \\\*\\\)D\\\.\[0-9\]+\\\]" "gimple" } }
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 0674a27501b..b8769d5a071 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -452,10 +452,12 @@ dump_omp_iterators (pretty_printer *pp, tree iter, int
spc, dump_flags_t flags)
{
pp_string (pp, ", loop_label=");
dump_generic_node (pp, TREE_VEC_ELT (iter, 6), spc, flags, false);
- pp_string (pp, ", elems=");
- dump_generic_node (pp, TREE_VEC_ELT (iter, 7), spc, flags, false);
pp_string (pp, ", index=");
+ dump_generic_node (pp, TREE_VEC_ELT (iter, 7), spc, flags, false);
+ pp_string (pp, ", elems=");
dump_generic_node (pp, TREE_VEC_ELT (iter, 8), spc, flags, false);
+ pp_string (pp, ", elems_count=");
+ dump_generic_node (pp, TREE_VEC_ELT (iter, 9), spc, flags, false);
}
pp_right_paren (pp);
}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1c.c
b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1c.c
new file mode 100644
index 00000000000..4fbd292cd0d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1c.c
@@ -0,0 +1,58 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* Like target-map-iterators-1.c, but with non-constant iterator
+ start/end/step. */
+
+#include <stdlib.h>
+
+#define DIM1 300
+#define DIM2 15
+
+int mkarray (int *x[], int a, int b, int c)
+{
+ int expected = 0;
+
+ for (int i = a; i < b; i++)
+ {
+ if ((i + 1) % c == 0)
+ {
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = rand ();
+ expected += x[i][j];
+ }
+ }
+ else
+ x[i] = NULL;
+ }
+
+ return expected;
+}
+
+int doit (int a, int b, int c)
+{
+ int *x[DIM1];
+ int y;
+
+ int expected = mkarray (x, a, b, c);
+
+ #pragma omp target enter data map(to: x)
+ #pragma omp target map(iterator(i=a:b:c), to: x[i][:DIM2]) \
+ map(from: y)
+ {
+ y = 0;
+ for (int i = a; i < b; i++)
+ if ((i + 1) % c == 0)
+ for (int j = 0; j < DIM2; j++)
+ y += x[i][j];
+ }
+
+ return y - expected;
+}
+
+int main (void)
+{
+ return doit (2, 300, 3);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1d.c
b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1d.c
new file mode 100644
index 00000000000..452c7f21d08
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1d.c
@@ -0,0 +1,69 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* Like target-map-iterators-1.c, but with non-constant iterator
+ start/end/step in target enter/exit data. */
+
+#include <stdlib.h>
+
+#define DIM1 300
+#define DIM2 15
+
+int mkarray (int *x[], int a, int b, int c)
+{
+ int expected = 0;
+
+ for (int i = a; i < b; i++)
+ {
+ if ((i + 1) % c == 0)
+ {
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = rand ();
+ expected += x[i][j];
+ }
+ }
+ else
+ x[i] = NULL;
+ }
+
+ return expected;
+}
+
+int doit (int a, int b, int c)
+{
+ int *x[DIM1];
+ int y;
+
+ int expected = mkarray (x, a, b, c);
+
+ #pragma omp target enter data map(iterator(i=a:b:c), to: x[i][:DIM2]) \
+ map(to: x)
+ #pragma omp target map(from: y)
+ {
+ y = 0;
+ for (int i = a; i < b; i++)
+ if ((i + 1) % c == 0)
+ for (int j = 0; j < DIM2; j++)
+ {
+ y += x[i][j];
+ x[i][j] = 0;
+ }
+ }
+
+ #pragma omp target exit data map(iterator(i=a:b:c), from: x[i][:DIM2])
+ #pragma omp target exit data map(release: x)
+
+ for (int i = a; i < b; i++)
+ if ((i + 1) % c == 0)
+ for (int j = 0; j < DIM2; j++)
+ if (x[i][j] != 0)
+ return 1;
+ return y - expected;
+}
+
+int main (void)
+{
+ return doit (2, 300, 3);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1e.c
b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1e.c
new file mode 100644
index 00000000000..3259818fdf3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1e.c
@@ -0,0 +1,66 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* Like target-map-iterators-1.c, but with non-constant iterator
+ start/end/step in target data. */
+
+#include <stdlib.h>
+
+#define DIM1 300
+#define DIM2 15
+
+int mkarray (int *x[], int a, int b, int c)
+{
+ int expected = 0;
+
+ for (int i = a; i < b; i++)
+ {
+ if ((i + 1) % c == 0)
+ {
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = rand ();
+ expected += x[i][j];
+ }
+ }
+ else
+ x[i] = NULL;
+ }
+
+ return expected;
+}
+
+int doit (int a, int b, int c)
+{
+ int *x[DIM1];
+ int y;
+
+ int expected = mkarray (x, a, b, c);
+
+ #pragma omp target data map(iterator(i=a:b:c), tofrom: x[i][:DIM2]) \
+ map(to: x)
+ #pragma omp target map(from: y)
+ {
+ y = 0;
+ for (int i = a; i < b; i++)
+ if ((i + 1) % c == 0)
+ for (int j = 0; j < DIM2; j++)
+ {
+ y += x[i][j];
+ x[i][j] = 0;
+ }
+ }
+
+ for (int i = a; i < b; i++)
+ if ((i + 1) % c == 0)
+ for (int j = 0; j < DIM2; j++)
+ if (x[i][j] != 0)
+ return 1;
+ return y - expected;
+}
+
+int main (void)
+{
+ return doit (2, 300, 3);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1f.c
b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1f.c
new file mode 100644
index 00000000000..353e596c5d8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1f.c
@@ -0,0 +1,71 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* Like target-map-iterators-1.c, but with non-constant iterator
+ start/end/step in target update. */
+
+#include <stdlib.h>
+
+#define DIM1 300
+#define DIM2 15
+
+void mkarray (int *x[], int a, int b, int c)
+{
+ for (int i = a; i < b; i++)
+ {
+ if ((i + 1) % c == 0)
+ {
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+ for (int j = 0; j < DIM2; j++)
+ x[i][j] = 0;
+ }
+ else
+ x[i] = NULL;
+ }
+}
+
+int doit (int a, int b, int c)
+{
+ int *x[DIM1];
+ int y;
+
+ /* Initialize the arrays to zero. */
+ mkarray (x, a, b, c);
+
+ #pragma omp target enter data map(iterator(i=a:b:c), to: x[i][:DIM2]) \
+ map(to: x)
+
+ /* Fill in some "real" values for the array on the host side. */
+ int expected = 0;
+ for (int i = a; i < b; i++)
+ {
+ if ((i + 1) % c == 0)
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = rand ();
+ expected += x[i][j];
+ }
+ }
+
+ /* Push the "real" values to the target. */
+ #pragma omp target update to(iterator(i=a:b:c): x[i][:DIM2])
+
+ #pragma omp target map(from: y)
+ {
+ y = 0;
+ for (int i = a; i < b; i++)
+ if ((i + 1) % c == 0)
+ for (int j = 0; j < DIM2; j++)
+ y += x[i][j];
+ }
+
+ #pragma omp target exit data map(iterator(i=a:b:c), from: x[i][:DIM2])
+ #pragma omp target exit data map(release: x)
+
+ return y - expected;
+}
+
+int main (void)
+{
+ return doit (2, 300, 3);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-4.c
b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-4.c
new file mode 100644
index 00000000000..621736762d4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-4.c
@@ -0,0 +1,48 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* Test transfer of dynamically-allocated arrays to target using map
+ iterators with non-constant bounds. */
+
+#include <stdlib.h>
+
+#define DIM1 8
+#define DIM2 15
+
+int mkarray (int *x[], int *dim1)
+{
+ int expected = 0;
+ *dim1 = DIM1;
+ for (int i = 0; i < DIM1; i++)
+ {
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = rand ();
+ expected += x[i][j];
+ }
+ }
+
+ return expected;
+}
+
+int main (void)
+{
+ int *x[DIM1];
+ int y;
+ int dim1;
+
+ int expected = mkarray (x, &dim1);
+
+ #pragma omp target enter data map(to: x)
+ #pragma omp target map(iterator(i=0:dim1), to: x[i][:DIM2]) \
+ map(from: y)
+ {
+ y = 0;
+ for (int i = 0; i < dim1; i++)
+ for (int j = 0; j < DIM2; j++)
+ y += x[i][j];
+ }
+
+ return y - expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-5.c
b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-5.c
new file mode 100644
index 00000000000..54b481877ed
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-5.c
@@ -0,0 +1,59 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* Test transfer of dynamically-allocated arrays to target using map
+ iterators, with multiple iterators, function calls and non-constant
+ bounds in the iterator expression. */
+
+#include <stdlib.h>
+
+#define DIM1 16
+#define DIM2 15
+
+int mkarrays (int *x[], int *y[], int *dim1)
+{
+ int expected = 0;
+
+ *dim1 = DIM1;
+ for (int i = 0; i < DIM1; i++)
+ {
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+ y[i] = (int *) malloc (sizeof (int));
+ *y[i] = rand ();
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = rand ();
+ expected += x[i][j] * *y[i];
+ }
+ }
+
+ return expected;
+}
+
+int f (int i, int j)
+{
+ return i * 4 + j;
+}
+
+int main (void)
+{
+ int *x[DIM1], *y[DIM1];
+ int sum;
+
+ int dim1;
+ int expected = mkarrays (x, y, &dim1);
+ int dim1_4 = dim1 / 4;
+
+ #pragma omp target enter data map(to: x, y)
+ #pragma omp target map(iterator(i=0:dim1_4, j=0:4), to: x[f(i, j)][:DIM2]) \
+ map(iterator(i=0:dim1), to: y[i][:1]) \
+ map(from: sum)
+ {
+ sum = 0;
+ for (int i = 0; i < dim1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j] * y[i][0];
+ }
+
+ return sum - expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-4.c
b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-4.c
new file mode 100644
index 00000000000..810b88165ce
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-4.c
@@ -0,0 +1,66 @@
+/* { dg-do run } */
+
+/* Test target enter data and target update to the target using map
+ iterators with non-constant bounds. */
+
+#include <stdlib.h>
+
+#define DIM1 8
+#define DIM2 15
+
+int mkarray (int *x[], int *dim1)
+{
+ int expected = 0;
+ *dim1 = DIM1;
+ for (int i = 0; i < DIM1; i++)
+ {
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = rand ();
+ expected += x[i][j];
+ }
+ }
+
+ return expected;
+}
+
+int main (void)
+{
+ int *x[DIM1];
+ int sum, dim1;
+ int expected = mkarray (x, &dim1);
+
+ #pragma omp target enter data map(to: x[:DIM1])
+ #pragma omp target enter data map(iterator(i=0:dim1), to: x[i][:DIM2])
+ #pragma omp target map(from: sum)
+ {
+ sum = 0;
+ for (int i = 0; i < dim1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j];
+ }
+
+ if (sum != expected)
+ return 1;
+
+ expected = 0;
+ for (int i = 0; i < dim1; i++)
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] *= rand ();
+ expected += x[i][j];
+ }
+
+ #pragma omp target update to(iterator(i=0:dim1): x[i][:DIM2])
+
+ #pragma omp target map(from: sum)
+ {
+ sum = 0;
+ for (int i = 0; i < dim1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j];
+ }
+
+ return sum != expected;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-map-iterators-1c.f90
b/libgomp/testsuite/libgomp.fortran/target-map-iterators-1c.f90
new file mode 100644
index 00000000000..d38566ba815
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-map-iterators-1c.f90
@@ -0,0 +1,58 @@
+! { dg-do run }
+
+! Like target-map-iterators-1.f90, but with non-constant iterator
+! start/end/step.
+
+program test
+ implicit none
+
+ integer, parameter :: DIM1 = 30
+ integer, parameter :: DIM2 = 15
+
+ type :: array_ptr
+ integer, pointer :: arr(:)
+ end type
+
+ type (array_ptr) :: x(DIM1)
+ integer :: expected
+
+ expected = mkarray (2, 29, 3)
+ if (doit (2, 29, 3) .ne. expected) stop 1
+
+contains
+ integer function mkarray (a, b, c)
+ integer :: a, b, c
+ integer :: exp = 0
+ integer :: i, j
+
+ do i = a, b
+ if (mod(i+1, c) == 0) then
+ allocate (x(i)%arr(DIM2))
+ do j = 1, DIM2
+ x(i)%arr(j) = i * j
+ exp = exp + x(i)%arr(j)
+ end do
+ end if
+ end do
+
+ mkarray = exp
+ end function
+
+ integer function doit (a, b, c)
+ integer :: a, b, c
+
+ integer:: sum, i, j
+ !$omp target map(iterator(i=a:b:c), to: x(i)%arr(:)) map(from: sum)
+ sum = 0
+ do i = a, b
+ if (mod(i+1, c) == 0) then
+ do j = 1, DIM2
+ sum = sum + x(i)%arr(j)
+ end do
+ end if
+ end do
+ !$omp end target
+
+ doit = sum
+ end function
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/target-map-iterators-4.f90
b/libgomp/testsuite/libgomp.fortran/target-map-iterators-4.f90
new file mode 100644
index 00000000000..85f6287bfcb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-map-iterators-4.f90
@@ -0,0 +1,48 @@
+! { dg-do run }
+
+! Test transfer of dynamically-allocated arrays to target using map
+! iterators with variable bounds.
+
+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
+ integer :: i_ubound
+
+ expected = mkarray (i_ubound)
+
+ !$omp target map(iterator(i=1:i_ubound), to: x(i)%arr(:)) map(from: sum)
+ sum = 0
+ do i = 1, i_ubound
+ 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 (ubound)
+ integer, intent(out) :: ubound
+ 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
+
+ ubound = DIM1
+ mkarray = exp
+ end function
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/target-map-iterators-5.f90
b/libgomp/testsuite/libgomp.fortran/target-map-iterators-5.f90
new file mode 100644
index 00000000000..9fb7e52279f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-map-iterators-5.f90
@@ -0,0 +1,55 @@
+! { dg-do run }
+
+! Test transfer of dynamically-allocated arrays to target using map
+! iterators, with multiple iterators, function calls and non-constant
+! bounds 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
+ integer :: i_ubound
+ integer :: k_ubound
+
+ expected = mkarrays (k_ubound)
+ i_ubound = k_ubound / 4 - 1
+
+ !$omp target map(iterator(i=0:i_ubound, j=0:3), to: x(i * 4 + j + 1)%arr(:))
&
+ !$omp map(iterator(k=1:k_ubound), 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 (ubound)
+ integer, intent(out) :: ubound
+ 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
+
+ ubound = DIM1
+ mkarrays = exp
+ end function
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/target-update-iterators-4.f90
b/libgomp/testsuite/libgomp.fortran/target-update-iterators-4.f90
new file mode 100644
index 00000000000..9f138aa6a5e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-update-iterators-4.f90
@@ -0,0 +1,70 @@
+! { dg-do run }
+
+! Test target enter data and target update to the target using map
+! iterators with non-constant bounds.
+
+program test
+ 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, ubound
+
+ expected = mkarray (x, ubound)
+
+ !$omp target enter data map(to: x)
+ !$omp target enter data map(iterator(i=1:ubound), to: x(i)%arr(:))
+ !$omp target map(from: sum)
+ sum = 0
+ do i = 1, ubound
+ do j = 1, DIM2
+ sum = sum + x(i)%arr(j)
+ end do
+ end do
+ !$omp end target
+
+ print *, sum, expected
+ if (sum .ne. expected) stop 1
+
+ expected = 0
+ do i = 1, ubound
+ do j = 1, DIM2
+ x(i)%arr(j) = x(i)%arr(j) * i * j
+ expected = expected + x(i)%arr(j)
+ end do
+ end do
+
+ !$omp target update to(iterator(i=1:ubound): x(i)%arr(:))
+
+ !$omp target map(from: sum)
+ sum = 0
+ do i = 1, ubound
+ do j = 1, DIM2
+ sum = sum + x(i)%arr(j)
+ end do
+ end do
+ !$omp end target
+
+ if (sum .ne. expected) stop 2
+contains
+ integer function mkarray (x, bound)
+ type (array_ptr), intent(inout) :: x(DIM1)
+ integer, intent(out) :: bound
+ 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
+
+ bound = DIM1
+ mkarray = exp
+ end function
+end program
--
2.39.5
