This patch builds on the previous patch series implementing OpenMP
iterators at:
https://gcc.gnu.org/pipermail/gcc-patches/2024-November/670333.html
This patch removes the limitation that the lower and upper bounds and
strides of the iterator must be compile-time constants - they can now be
anything that results in an integer expression.
This means that the internal arrays used to hold the expanded clause
expressions must now be dynamically allocated instead of statically
(static arrays are still used if the iteration count is determined to be
a compile-time constant). Calls to malloc the arrays are generated just
before the loops created to expand the iterators, and calls to free them
are generated after the target statement. The malloc calls are added in
the omplower stage rather than when the loops are being built during
Gimplification because clauses can still get moved around and removed at
that point.
Kwok
From d277a9539ad78fae9eb97a156949320623060a0d 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] 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.
2024-12-10 Kwok Cheung Yeung <kcye...@baylibre.com>
gcc/
* gimplify.cc (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 array length.
* 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 | 34 +++---
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 ++++++++++++
15 files changed, 509 insertions(+), 44 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 b761082f61f..ea0aa79e3d9 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9442,25 +9442,20 @@ 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);
- TREE_ADDRESSABLE (elems) = 1;
- gimple_add_tmp_var (elems);
- }
- else
- {
- /* Handle dynamic sizes. */
- sorry ("dynamic iterator sizes not implemented yet");
- }
+ 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);
+ 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);
@@ -9468,10 +9463,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) = arr_length;
TREE_CHAIN (new_iterator) = TREE_CHAIN (OMP_CLAUSE_ITERATORS (c));
OMP_CLAUSE_ITERATORS (c) = new_iterator;
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 1db336fd7ba..976357e835d 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -12649,20 +12649,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.
@@ -12676,15 +12686,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);
@@ -12693,6 +12713,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 = 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. */
@@ -14378,9 +14454,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 23b21a46eb2..107cac0870b 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=D\\\.\[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=D\\\.\[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=D\\\.\[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=D\\\.\[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=D\\\.\[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=D\\\.\[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=D\\\.\[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=D\\\.\[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..5c33ab93f25
--- /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 \\(D\\\.\[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 21b5aa88083..ad8e77c235f 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=D\.\[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=D\.\[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=D\.\[0-9\]+,
elems_count=401\\):\\*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=D\.\[0-9\]+,
elems_count=21\\):\\*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 abb60f32aab..02c821daafe 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=D\\\.\[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=D\\\.\[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=D\\\.\[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=D\\\.\[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=D\\\.\[0-9\]+,
elems_count=35\\):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=D\\\.\[0-9\]+,
elems_count=55\\):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=D\\\.\[0-9\]+,
elems_count=35\\):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=D\\\.\[0-9\]+,
elems_count=55\\):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..6e2a4558ce6
--- /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 \\(D\\\.\[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 7629d95aa2b..c91a2ff0b18 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=D\\\.\[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=D\\\.\[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=D\\\.\[0-9\]+, elems_count=1327\\):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=D\\\.\[0-9\]+,
elems_count=35\\):MEM <\[^>\]+> \\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]"
"gimple" } }
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 0fd7c457308..7778e9abe94 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
