https://gcc.gnu.org/g:78879fbac86d326753aa73984df1771171e35f27
commit 78879fbac86d326753aa73984df1771171e35f27 Author: Kwok Cheung Yeung <[email protected]> Date: Sun Jun 7 01:37:17 2026 +0000 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-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]> Diff: --- gcc/gimplify.cc | 42 +++++----- gcc/gimplify.h | 1 + gcc/omp-low.cc | 93 ++++++++++++++++++++-- .../c-c++-common/gomp/target-map-iterators-3.c | 8 +- .../c-c++-common/gomp/target-map-iterators-5.c | 14 ++++ .../c-c++-common/gomp/target-update-iterators-3.c | 4 +- .../gfortran.dg/gomp/target-map-iterators-3.f90 | 8 +- .../gfortran.dg/gomp/target-map-iterators-5.f90 | 21 +++++ .../gfortran.dg/gomp/target-map-iterators-6.f90 | 6 +- .../gfortran.dg/gomp/target-update-iterators-3.f90 | 6 +- gcc/tree-pretty-print.cc | 7 +- .../libgomp.c-c++-common/target-map-iterators-1c.c | 58 ++++++++++++++ .../libgomp.c-c++-common/target-map-iterators-1d.c | 69 ++++++++++++++++ .../libgomp.c-c++-common/target-map-iterators-1e.c | 66 +++++++++++++++ .../libgomp.c-c++-common/target-map-iterators-1f.c | 71 +++++++++++++++++ .../libgomp.c-c++-common/target-map-iterators-4.c | 48 +++++++++++ .../libgomp.c-c++-common/target-map-iterators-5.c | 59 ++++++++++++++ .../target-update-iterators-4.c | 66 +++++++++++++++ .../libgomp.fortran/target-map-iterators-1c.f90 | 58 ++++++++++++++ .../libgomp.fortran/target-map-iterators-4.f90 | 48 +++++++++++ .../libgomp.fortran/target-map-iterators-5.f90 | 55 +++++++++++++ .../libgomp.fortran/target-update-iterators-4.f90 | 70 ++++++++++++++++ 22 files changed, 835 insertions(+), 43 deletions(-) diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index f2d9b70d61b8..c23792d06281 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -10215,6 +10215,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 @@ -10273,27 +10284,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); @@ -10303,6 +10308,7 @@ build_omp_iterators_loops (tree *list_p, gimple_seq *loops_seq_p) OMP_ITERATOR_LABEL (new_iterator) = loop.body_label; OMP_ITERATOR_INDEX (new_iterator) = loop.index; OMP_ITERATOR_ELEMS (new_iterator) = elems; + OMP_ITERATOR_COUNT (new_iterator) = 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 caa35b426bd6..2a095e85f5f8 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 a690190accd2..db29566946da 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -12775,14 +12775,24 @@ lower_omp_map_iterator_expr (tree expr, tree c, gomp_target *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 @@ -12803,9 +12813,19 @@ lower_omp_map_iterator_size (tree size, tree c, gomp_target *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); @@ -12814,6 +12834,63 @@ 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 = OMP_ITERATOR_ELEMS (iters); + if (!POINTER_TYPE_P (TREE_TYPE (elems))) + continue; + tree arr_length + = omp_iterator_elems_length (OMP_ITERATOR_COUNT (iters)); + 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 = OMP_ITERATOR_INDEX (iters); + 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 = OMP_ITERATOR_ELEMS (OMP_CLAUSE_ITERATORS (c)); + 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. */ @@ -14752,9 +14829,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 c31336538deb..82927edbbf86 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 000000000000..58daa0e73bf1 --- /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 0057cf76d87d..00da5138e77e 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 7dad2a69c611..1099955c2cc8 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 000000000000..2de312a8b9ae --- /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 b83481920b60..5cc36f0d024b 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 d9c92cf46790..a8dffcf09ba0 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 b470a9dee7b5..e64c3a7f8414 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -452,11 +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, OMP_ITERATOR_LABEL (iter), spc, flags, false); - pp_string (pp, ", elems="); - dump_generic_node (pp, OMP_ITERATOR_ELEMS (iter), spc, flags, false); pp_string (pp, ", index="); dump_generic_node (pp, OMP_ITERATOR_INDEX (iter), spc, flags, false); - /* The count field is not used yet. */ + pp_string (pp, ", elems="); + dump_generic_node (pp, OMP_ITERATOR_ELEMS (iter), spc, flags, false); + pp_string (pp, ", elems_count="); + dump_generic_node (pp, OMP_ITERATOR_COUNT (iter), 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 000000000000..4fbd292cd0da --- /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 000000000000..452c7f21d083 --- /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 000000000000..3259818fdf39 --- /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 000000000000..353e596c5d81 --- /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 000000000000..621736762d40 --- /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 000000000000..54b481877ed8 --- /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 000000000000..810b88165ce0 --- /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 000000000000..d38566ba815d --- /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 000000000000..85f6287bfcba --- /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 000000000000..9fb7e52279f4 --- /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 000000000000..9f138aa6a5e3 --- /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
