+2025-04-17 Kwok Cheung Yeung <kcyeung@baylibre.com>
+
+ * 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.
+
2025-04-17 Kwok Cheung Yeung <kcyeung@baylibre.com>
* gimplify.cc (compute_omp_iterator_count): Account for difference
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
}
/* 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);
/* 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;
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);
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.
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);
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. */
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);
+2025-04-17 Kwok Cheung Yeung <kcyeung@baylibre.com>
+
+ * 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.
+
2025-04-17 Kwok Cheung Yeung <kcyeung@baylibre.com>
* gfortran.dg/gomp/target-update-iterators-1.f90: New.
/* { 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" } } */
--- /dev/null
+/* { 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" } } */
/* { 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" } } */
! { 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" } }
--- /dev/null
+! { 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" } }
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" } }
{
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);
}
+2025-04-17 Kwok Cheung Yeung <kcyeung@baylibre.com>
+
+ * 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.
+
2025-04-17 Kwok Cheung Yeung <kcyeung@baylibre.com>
* testsuite/libgomp.fortran/target-update-iterators-1.f90: New.
--- /dev/null
+/* { 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;
+}
--- /dev/null
+/* { 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;
+}
--- /dev/null
+/* { 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;
+}
--- /dev/null
+! { 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
--- /dev/null
+! { 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
--- /dev/null
+! { 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
projects / thirdparty / gcc.git / commitdiff
