The previous version was posted at:
https://gcc.gnu.org/pipermail/gcc-patches/2024-December/671630.html
This patch is largely the same as the previous version, but calculation
of the size of the iterator elements array is separated into another
function for reuse later.From 607ceefb93136fd14471d0a9383464776228f074 Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <kcye...@baylibre.com>
Date: Thu, 12 Dec 2024 21:22:20 +0000
Subject: [PATCH 06/11] openmp: Add support for non-constant iterator
parameters in map, to and from clauses
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-update-iterators-3.f90: Update expected
Gimple output.
libgomp/
* 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-4.f90: New.
* testsuite/libgomp.fortran/target-map-iterators-5.f90: New.
* testsuite/libgomp.fortran/target-update-iterators-4.f90: New.
---
gcc/gimplify.cc | 44 ++++----
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-update-iterators-3.f90 | 6 +-
gcc/tree-pretty-print.cc | 6 +-
.../target-map-iterators-4.c | 48 +++++++++
.../target-map-iterators-5.c | 59 +++++++++++
.../target-update-iterators-4.c | 66 ++++++++++++
.../target-map-iterators-4.f90 | 48 +++++++++
.../target-map-iterators-5.f90 | 61 +++++++++++
.../target-update-iterators-4.f90 | 70 ++++++++++++
16 files changed, 518 insertions(+), 46 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-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-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 52d0e11de8b..f8f649c7154 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9973,6 +9973,13 @@ struct iterator_loop_info_t
typedef hash_map<tree, iterator_loop_info_t> iterator_loop_info_map_t;
+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
@@ -10028,27 +10035,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);
@@ -10056,10 +10057,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 80c335e7c2c..ca970cff786 100644
--- a/gcc/gimplify.h
+++ b/gcc/gimplify.h
@@ -79,6 +79,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 01a33c68ac7..5b1b0fe256b 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -12662,20 +12662,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.
@@ -12689,15 +12699,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);
@@ -12706,6 +12726,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. */
@@ -14391,9 +14467,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 62df42ffde1..87b32e4ad77 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..e79e6a587da
--- /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 ef55216876f..6618962de41 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..f620d1c1517
--- /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\]+\\);" 3 "omplower" } }
+! { dg-final { scan-tree-dump-times "__builtin_free
\\(omp_iter_data\\\.\[0-9\]+\\);" 3 "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 6ccb52d8eaf..fa4075fdbe1 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-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-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..4c47ee57883
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-map-iterators-5.f90
@@ -0,0 +1,61 @@
+! { 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(f (i, j))%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
+
+ integer function f (i, j)
+ integer, intent(in) :: i, j
+
+ f = i * 4 + j + 1
+ 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.43.0
