v1: https://gcc.gnu.org/pipermail/gcc-patches/2024-May/652682.html
v2: https://gcc.gnu.org/pipermail/gcc-patches/2024-September/662140.html
v3: https://gcc.gnu.org/pipermail/gcc-patches/2024-October/664543.html
v4: https://gcc.gnu.org/pipermail/gcc-patches/2024-November/670335.html
This patch is largely unchanged from v4.
gimple_omp_target_iterator_loops_ptr now takes a gimple* rather than a
gomp_target* to avoid a couple of as_a casts in various places.From d3ad9940694b33a5428dbd5e52f46ec4da799419 Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <kcye...@baylibre.com>
Date: Sat, 3 May 2025 20:24:26 +0000
Subject: [PATCH 02/11] openmp: Add support for iterators in map clauses
(C/C++)
This adds preliminary support for iterators in map clauses within OpenMP
'target' constructs (which includes constructs such as 'target enter data').
Iterators with non-constant loop bounds are not currently supported.
gcc/c/
* c-parser.cc (c_parser_omp_clause_map): Parse 'iterator' modifier.
* c-typeck.cc (c_finish_omp_clauses): Finish iterators. Apply
iterators to generated clauses.
gcc/cp/
* parser.cc (cp_parser_omp_clause_map): Parse 'iterator' modifier.
* semantics.cc (finish_omp_clauses): Finish iterators. Apply
iterators to generated clauses.
gcc/
* gimple-pretty-print.cc (dump_gimple_omp_target): Print expanded
iterator loops.
* gimple.cc (gimple_build_omp_target): Add argument for iterator
loops sequence. Initialize iterator loops field.
* gimple.def (GIMPLE_OMP_TARGET): Set GSS symbol to GSS_OMP_TARGET.
* gimple.h (gomp_target): Set GSS symbol to GSS_OMP_TARGET. Add extra
field for iterator loops.
(gimple_build_omp_target): Add argument for iterator loops sequence.
(gimple_omp_target_iterator_loops): New.
(gimple_omp_target_iterator_loops_ptr): New.
(gimple_omp_target_set_iterator_loops): New.
* gimplify.cc (find_var_decl): New.
(copy_omp_iterator): New.
(remap_omp_iterator_var_1): New.
(remap_omp_iterator_var): New.
(remove_unused_omp_iterator_vars): New.
(struct iterator_loop_info_t): New type.
(iterator_loop_info_map_t): New type.
(build_omp_iterators_loops): New.
(enter_omp_iterator_loop_context_1): New.
(enter_omp_iterator_loop_context): New.
(enter_omp_iterator_loop_context): New.
(exit_omp_iterator_loop_context): New.
(gimplify_adjust_omp_clauses): Add argument for iterator loop
sequence. Gimplify the clause decl and size into the iterator
loop if iterators are used.
(gimplify_omp_workshare): Call remove_unused_omp_iterator_vars and
build_omp_iterators_loops for OpenMP target expressions. Add
loop sequence as argument when calling gimplify_adjust_omp_clauses
and building the Gimple statement.
* gimplify.h (enter_omp_iterator_loop_context): New prototype.
(exit_omp_iterator_loop_context): New prototype.
* gsstruct.def (GSS_OMP_TARGET): New.
* omp-low.cc (lower_omp_map_iterator_expr): New.
(lower_omp_map_iterator_size): New.
(finish_omp_map_iterators): New.
(lower_omp_target): Add sorry if iterators used with deep mapping.
Call lower_omp_map_iterator_expr before assigning to sender ref.
Call lower_omp_map_iterator_size before setting the size. Insert
iterator loop sequence before the statements for the target clause.
* tree-nested.cc (convert_nonlocal_reference_stmt): Walk the iterator
loop sequence of OpenMP target statements.
(convert_local_reference_stmt): Likewise.
(convert_tramp_reference_stmt): Likewise.
* tree-pretty-print.cc (dump_omp_iterators): Dump extra iterator
information if present.
(dump_omp_clause): Call dump_omp_iterators for iterators in map
clauses.
* tree.cc (omp_clause_num_ops): Add operand for OMP_CLAUSE_MAP.
(walk_tree_1): Do not walk last operand of OMP_CLAUSE_MAP.
* tree.h (OMP_CLAUSE_HAS_ITERATORS): New.
(OMP_CLAUSE_ITERATORS): New.
gcc/testsuite/
* c-c++-common/gomp/map-6.c (foo): Amend expected error message.
* c-c++-common/gomp/target-map-iterators-1.c: New.
* c-c++-common/gomp/target-map-iterators-2.c: New.
* c-c++-common/gomp/target-map-iterators-3.c: New.
* c-c++-common/gomp/target-map-iterators-4.c: New.
libgomp/
* target.c (kind_to_name): New.
(gomp_merge_iterator_maps): New.
(gomp_map_vars_internal): Call gomp_merge_iterator_maps. Copy
address of only the first iteration to target vars. Free allocated
variables.
* testsuite/libgomp.c-c++-common/target-map-iterators-1.c: New.
* testsuite/libgomp.c-c++-common/target-map-iterators-2.c: New.
* testsuite/libgomp.c-c++-common/target-map-iterators-3.c: New.
Co-authored-by: Andrew Stubbs <a...@baylibre.com>
---
gcc/c/c-parser.cc | 54 ++-
gcc/c/c-typeck.cc | 20 +-
gcc/cp/parser.cc | 54 ++-
gcc/cp/semantics.cc | 20 +-
gcc/gimple-pretty-print.cc | 6 +
gcc/gimple.cc | 8 +-
gcc/gimple.def | 2 +-
gcc/gimple.h | 43 +-
gcc/gimplify.cc | 397 +++++++++++++++++-
gcc/gimplify.h | 4 +
gcc/gsstruct.def | 1 +
gcc/omp-low.cc | 84 +++-
gcc/testsuite/c-c++-common/gomp/map-6.c | 20 +-
.../gomp/target-map-iterators-1.c | 23 +
.../gomp/target-map-iterators-2.c | 25 ++
.../gomp/target-map-iterators-3.c | 23 +
.../gomp/target-map-iterators-4.c | 18 +
gcc/tree-nested.cc | 8 +
gcc/tree-pretty-print.cc | 14 +
gcc/tree.cc | 5 +-
gcc/tree.h | 8 +
libgomp/target.c | 130 +++++-
.../target-map-iterators-1.c | 47 +++
.../target-map-iterators-2.c | 44 ++
.../target-map-iterators-3.c | 56 +++
25 files changed, 1057 insertions(+), 57 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-map-iterators-4.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 0c3e3e2889c..e60b3821481 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -20036,7 +20036,7 @@ c_parser_omp_clause_doacross (c_parser *parser, tree
list)
map ( [map-type-modifier[,] ...] map-kind: variable-list )
map-type-modifier:
- always | close */
+ always | close | present | iterator (iterators-definition) */
static tree
c_parser_omp_clause_map (c_parser *parser, tree list, bool declare_mapper_p)
@@ -20051,15 +20051,35 @@ c_parser_omp_clause_map (c_parser *parser, tree list,
bool declare_mapper_p)
int pos = 1;
int map_kind_pos = 0;
- while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)
+ int iterator_length = 0;
+ for (;;)
{
- if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON)
+ c_token *tok = c_parser_peek_nth_token_raw (parser, pos);
+ if (tok->type != CPP_NAME)
+ break;
+
+ const char *p = IDENTIFIER_POINTER (tok->value);
+ c_token *next_tok = c_parser_peek_nth_token_raw (parser, pos + 1);
+ if (strcmp (p, "iterator") == 0 && next_tok->type == CPP_OPEN_PAREN)
+ {
+ unsigned n = pos + 2;
+ if (c_parser_check_balanced_raw_token_sequence (parser, &n)
+ && c_parser_peek_nth_token_raw (parser, n)->type
+ == CPP_CLOSE_PAREN)
+ {
+ iterator_length = n - pos + 1;
+ pos = n;
+ next_tok = c_parser_peek_nth_token_raw (parser, pos + 1);
+ }
+ }
+
+ if (next_tok->type == CPP_COLON)
{
map_kind_pos = pos;
break;
}
- if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
+ if (next_tok->type == CPP_COMMA)
pos++;
else if (c_parser_peek_nth_token_raw (parser, pos + 1)->type
== CPP_OPEN_PAREN)
@@ -20081,6 +20101,7 @@ c_parser_omp_clause_map (c_parser *parser, tree list,
bool declare_mapper_p)
int present_modifier = 0;
int mapper_modifier = 0;
tree mapper_name = NULL_TREE;
+ tree iterators = NULL_TREE;
for (int pos = 1; pos < map_kind_pos; ++pos)
{
c_token *tok = c_parser_peek_token (parser);
@@ -20114,6 +20135,17 @@ c_parser_omp_clause_map (c_parser *parser, tree list,
bool declare_mapper_p)
close_modifier++;
c_parser_consume_token (parser);
}
+ else if (strcmp ("iterator", p) == 0)
+ {
+ if (iterators)
+ {
+ c_parser_error (parser, "too many %<iterator%> modifiers");
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ iterators = c_parser_omp_iterators (parser);
+ pos += iterator_length - 1;
+ }
else if (strcmp ("mapper", p) == 0)
{
c_parser_consume_token (parser);
@@ -20187,8 +20219,8 @@ c_parser_omp_clause_map (c_parser *parser, tree list,
bool declare_mapper_p)
else
{
c_parser_error (parser, "%<map%> clause with map-type modifier other "
- "than %<always%>, %<close%>, %<mapper%> or "
- "%<present%>");
+ "than %<always%>, %<close%>, %<iterator%>, "
+ "%<mapper%> or %<present%>");
parens.skip_until_found_close (parser);
return list;
}
@@ -20237,9 +20269,19 @@ c_parser_omp_clause_map (c_parser *parser, tree list,
bool declare_mapper_p)
tree last_new = NULL_TREE;
+ if (iterators)
+ {
+ tree block = pop_scope ();
+ if (iterators == error_mark_node)
+ iterators = NULL_TREE;
+ else
+ TREE_VEC_ELT (iterators, 5) = block;
+ }
+
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
{
OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ OMP_CLAUSE_ITERATORS (c) = iterators;
last_new = c;
}
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index d2a857de9a6..7f0d64dafd0 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -16186,7 +16186,14 @@ c_finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
/* We've reached the end of a list of expanded nodes. Reset the group
start pointer. */
if (c == grp_sentinel)
- grp_start_p = NULL;
+ {
+ if (grp_start_p
+ && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p))
+ for (tree gc = *grp_start_p; gc != grp_sentinel;
+ gc = OMP_CLAUSE_CHAIN (gc))
+ OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p);
+ grp_start_p = NULL;
+ }
switch (OMP_CLAUSE_CODE (c))
{
@@ -16941,6 +16948,12 @@ c_finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
remove = true;
break;
}
+ if (OMP_CLAUSE_ITERATORS (c)
+ && c_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c)))
+ {
+ t = error_mark_node;
+ break;
+ }
/* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
@@ -17672,6 +17685,11 @@ c_finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
pc = &OMP_CLAUSE_CHAIN (c);
}
+ if (grp_start_p
+ && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p))
+ for (tree gc = *grp_start_p; gc; gc = OMP_CLAUSE_CHAIN (gc))
+ OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p);
+
if (simdlen
&& safelen
&& tree_int_cst_lt (OMP_CLAUSE_SAFELEN_EXPR (safelen),
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 44a78324c6e..b038332e732 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -42479,16 +42479,34 @@ cp_parser_omp_clause_map (cp_parser *parser, tree
list, bool declare_mapper_p)
int pos = 1;
int map_kind_pos = 0;
- while (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_NAME
- || cp_lexer_peek_nth_token (parser->lexer, pos)->keyword == RID_DELETE)
+ int iterator_length = 0;
+ for (;;)
{
- if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COLON)
+ cp_token *tok = cp_lexer_peek_nth_token (parser->lexer, pos);
+ if (!(tok->type == CPP_NAME || tok->keyword == RID_DELETE))
+ break;
+
+ cp_token *next_tok = cp_lexer_peek_nth_token (parser->lexer, pos + 1);
+ if (tok->type == CPP_NAME
+ && strcmp (IDENTIFIER_POINTER (tok->u.value), "iterator") == 0
+ && next_tok->type == CPP_OPEN_PAREN)
+ {
+ int n = cp_parser_skip_balanced_tokens (parser, pos + 1);
+ if (n != pos + 1)
+ {
+ iterator_length = n - pos;
+ pos = n - 1;
+ next_tok = cp_lexer_peek_nth_token (parser->lexer, n);
+ }
+ }
+
+ if (next_tok->type == CPP_COLON)
{
map_kind_pos = pos;
break;
}
- if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)
+ if (next_tok->type == CPP_COMMA)
pos++;
else if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type
== CPP_OPEN_PAREN)
@@ -42501,6 +42519,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list,
bool declare_mapper_p)
bool present_modifier = false;
bool mapper_modifier = false;
tree mapper_name = NULL_TREE;
+ tree iterators = NULL_TREE;
for (int pos = 1; pos < map_kind_pos; ++pos)
{
cp_token *tok = cp_lexer_peek_token (parser->lexer);
@@ -42539,6 +42558,21 @@ cp_parser_omp_clause_map (cp_parser *parser, tree
list, bool declare_mapper_p)
close_modifier = true;
cp_lexer_consume_token (parser->lexer);
}
+ else if (strcmp ("iterator", p) == 0)
+ {
+ if (iterators)
+ {
+ cp_parser_error (parser, "too many %<iterator%> modifiers");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+ begin_scope (sk_omp, NULL);
+ iterators = cp_parser_omp_iterators (parser);
+ pos += iterator_length - 1;
+ }
else if (strcmp ("mapper", p) == 0)
{
cp_lexer_consume_token (parser->lexer);
@@ -42627,7 +42661,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list,
bool declare_mapper_p)
{
cp_parser_error (parser, "%<map%> clause with map-type modifier "
"other than %<always%>, %<close%>, "
- "%<mapper%> or %<present%>");
+ "%<iterator%>, %<mapper%> or %<present%>");
cp_parser_skip_to_closing_parenthesis (parser,
/*recovering=*/true,
/*or_comma=*/false,
@@ -42691,9 +42725,19 @@ cp_parser_omp_clause_map (cp_parser *parser, tree
list, bool declare_mapper_p)
tree last_new = NULL_TREE;
+ if (iterators)
+ {
+ tree block = poplevel (1, 1, 0);
+ if (iterators == error_mark_node)
+ iterators = NULL_TREE;
+ else
+ TREE_VEC_ELT (iterators, 5) = block;
+ }
+
for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
{
OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ OMP_CLAUSE_ITERATORS (c) = iterators;
last_new = c;
}
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 07a601eb172..4f88d0702ce 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -7736,7 +7736,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type
ort)
/* We've reached the end of a list of expanded nodes. Reset the group
start pointer. */
if (c == grp_sentinel)
- grp_start_p = NULL;
+ {
+ if (grp_start_p
+ && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p))
+ for (tree gc = *grp_start_p; gc != grp_sentinel;
+ gc = OMP_CLAUSE_CHAIN (gc))
+ OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p);
+ grp_start_p = NULL;
+ }
switch (OMP_CLAUSE_CODE (c))
{
@@ -8985,6 +8992,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type
ort)
remove = true;
break;
}
+ if (OMP_CLAUSE_ITERATORS (c)
+ && cp_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c)))
+ {
+ t = error_mark_node;
+ break;
+ }
/* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
@@ -9891,6 +9904,11 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type
ort)
pc = &OMP_CLAUSE_CHAIN (c);
}
+ if (grp_start_p
+ && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p))
+ for (tree gc = *grp_start_p; gc; gc = OMP_CLAUSE_CHAIN (gc))
+ OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p);
+
if (reduction_seen < 0 && (ordered_seen || schedule_seen))
reduction_seen = -2;
diff --git a/gcc/gimple-pretty-print.cc b/gcc/gimple-pretty-print.cc
index 4e20b4cc371..6929cd0bca1 100644
--- a/gcc/gimple-pretty-print.cc
+++ b/gcc/gimple-pretty-print.cc
@@ -1837,6 +1837,12 @@ dump_gimple_omp_target (pretty_printer *pp, const
gomp_target *gs,
default:
gcc_unreachable ();
}
+ if (gimple_omp_target_iterator_loops (gs))
+ {
+ pp_string (pp, "// Expanded iterator loops for #pragma omp target\n");
+ dump_gimple_seq (pp, gimple_omp_target_iterator_loops (gs), spc, flags);
+ pp_newline (pp);
+ }
if (flags & TDF_RAW)
{
dump_gimple_fmt (pp, spc, flags, "%G%s <%+BODY <%S>%nCLAUSES <", gs,
diff --git a/gcc/gimple.cc b/gcc/gimple.cc
index 41908d4e29a..102e21fe5e5 100644
--- a/gcc/gimple.cc
+++ b/gcc/gimple.cc
@@ -1295,10 +1295,13 @@ gimple_build_omp_interop (tree clauses)
BODY is the sequence of statements that will be executed.
KIND is the kind of the region.
- CLAUSES are any of the construct's clauses. */
+ CLAUSES are any of the construct's clauses.
+ ITERATOR_LOOPS is an optional sequence containing constructed loops
+ for OpenMP iterators. */
gomp_target *
-gimple_build_omp_target (gimple_seq body, int kind, tree clauses)
+gimple_build_omp_target (gimple_seq body, int kind, tree clauses,
+ gimple_seq iterator_loops)
{
gomp_target *p
= as_a <gomp_target *> (gimple_alloc (GIMPLE_OMP_TARGET, 0));
@@ -1306,6 +1309,7 @@ gimple_build_omp_target (gimple_seq body, int kind, tree
clauses)
gimple_omp_set_body (p, body);
gimple_omp_target_set_clauses (p, clauses);
gimple_omp_target_set_kind (p, kind);
+ gimple_omp_target_set_iterator_loops (p, iterator_loops);
return p;
}
diff --git a/gcc/gimple.def b/gcc/gimple.def
index 54248a80aa6..3e1e13ebb2c 100644
--- a/gcc/gimple.def
+++ b/gcc/gimple.def
@@ -393,7 +393,7 @@ DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single",
GSS_OMP_SINGLE_LAYOUT)
DATA_ARG is a vec of 3 local variables in the parent function
containing data to be mapped to CHILD_FN. This is used to
implement the MAP clauses. */
-DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_PARALLEL_LAYOUT)
+DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_TARGET)
/* GIMPLE_OMP_TEAMS <BODY, CLAUSES, CHILD_FN, DATA_ARG> represents
#pragma omp teams
diff --git a/gcc/gimple.h b/gcc/gimple.h
index 268884677e1..6c3ab2b6d23 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -682,11 +682,14 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
};
/* GIMPLE_OMP_TARGET */
-struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
+struct GTY((tag("GSS_OMP_TARGET")))
gomp_target : public gimple_statement_omp_parallel_layout
{
- /* No extra fields; adds invariant:
- stmt->code == GIMPLE_OMP_TARGET. */
+ /* [ WORD 1-10 ] : base class */
+
+ /* [ WORD 11 ]
+ Iterator loops. */
+ gimple_seq iterator_loops;
};
/* GIMPLE_OMP_TASK */
@@ -1607,7 +1610,7 @@ gomp_scan *gimple_build_omp_scan (gimple_seq, tree);
gomp_sections *gimple_build_omp_sections (gimple_seq, tree);
gimple *gimple_build_omp_sections_switch (void);
gomp_single *gimple_build_omp_single (gimple_seq, tree);
-gomp_target *gimple_build_omp_target (gimple_seq, int, tree);
+gomp_target *gimple_build_omp_target (gimple_seq, int, tree, gimple_seq =
NULL);
gomp_teams *gimple_build_omp_teams (gimple_seq, tree);
gomp_atomic_load *gimple_build_omp_atomic_load (tree, tree,
enum omp_memory_order);
@@ -6380,6 +6383,38 @@ gimple_omp_target_set_data_arg (gomp_target
*omp_target_stmt,
}
+/* Return the Gimple sequence used to store loops for OpenMP iterators used
+ by OMP_TARGET_STMT. */
+
+inline gimple_seq
+gimple_omp_target_iterator_loops (const gomp_target *omp_target_stmt)
+{
+ return omp_target_stmt->iterator_loops;
+}
+
+
+/* Return a pointer to the Gimple sequence used to store loops for OpenMP
+ iterators used by OMP_TARGET GS. */
+
+inline gimple_seq *
+gimple_omp_target_iterator_loops_ptr (gimple *gs)
+{
+ gomp_target *omp_target_stmt = as_a <gomp_target *> (gs);
+ return &omp_target_stmt->iterator_loops;
+}
+
+
+/* Set ITERATOR_LOOPS to be the Gimple sequence used to store loops
+ constructed for OpenMP iterators in OMP_TARGET_STMT. */
+
+inline void
+gimple_omp_target_set_iterator_loops (gomp_target *omp_target_stmt,
+ gimple_seq iterator_loops)
+{
+ omp_target_stmt->iterator_loops = iterator_loops;
+}
+
+
/* Return the clauses associated with OMP_TEAMS GS. */
inline tree
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 34c976fc17e..6e4c12d0ebf 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9780,6 +9780,367 @@ build_omp_iterator_loop (tree it, gimple_seq *pre_p,
tree *last_bind)
return p;
}
+
+/* Callback for walk_tree to find a VAR_DECL (stored in DATA) in the
+ tree TP. */
+
+static tree
+find_var_decl (tree *tp, int *, void *data)
+{
+ if (*tp == (tree) data)
+ return *tp;
+
+ return NULL_TREE;
+}
+
+/* Returns an element-by-element copy of OMP iterator tree IT. */
+
+static tree
+copy_omp_iterator (tree it, int elem_count = -1)
+{
+ if (elem_count < 0)
+ elem_count = TREE_VEC_LENGTH (it);
+ tree new_it = make_tree_vec (elem_count);
+ for (int i = 0; i < TREE_VEC_LENGTH (it); i++)
+ TREE_VEC_ELT (new_it, i) = TREE_VEC_ELT (it, i);
+
+ return new_it;
+}
+
+/* Helper function for walk_tree in remap_omp_iterator_var. */
+
+static tree
+remap_omp_iterator_var_1 (tree *tp, int *, void *data)
+{
+ tree old_var = ((tree *) data)[0];
+ tree new_var = ((tree *) data)[1];
+
+ if (*tp == old_var)
+ *tp = new_var;
+ return NULL_TREE;
+}
+
+/* Replace instances of OLD_VAR in TP with NEW_VAR. */
+
+static void
+remap_omp_iterator_var (tree *tp, tree old_var, tree new_var)
+{
+ tree vars[2] = { old_var, new_var };
+ walk_tree (tp, remap_omp_iterator_var_1, vars, NULL);
+}
+
+/* Scan through all clauses using OpenMP iterators in LIST_P. If any
+ clauses have iterators with variables that are not used by the clause
+ decl or size, issue a warning and replace the iterator with a copy with
+ the unused variables removed. */
+
+static void
+remove_unused_omp_iterator_vars (tree *list_p)
+{
+ auto_vec< vec<tree> > iter_vars;
+ auto_vec<tree> new_iterators;
+
+ for (tree c = *list_p; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (!OMP_CLAUSE_HAS_ITERATORS (c))
+ continue;
+ auto_vec<tree> vars;
+ bool need_new_iterators = false;
+ for (tree it = OMP_CLAUSE_ITERATORS (c); it; it = TREE_CHAIN (it))
+ {
+ tree var = TREE_VEC_ELT (it, 0);
+ tree t = walk_tree (&OMP_CLAUSE_DECL (c), find_var_decl, var, NULL);
+ if (t == NULL_TREE)
+ t = walk_tree (&OMP_CLAUSE_SIZE (c), find_var_decl, var, NULL);
+ if (t == NULL_TREE)
+ {
+ need_new_iterators = true;
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH)
+ warning_at (OMP_CLAUSE_LOCATION (c), 0,
+ "iterator variable %qE not used in clause "
+ "expression", DECL_NAME (var));
+ }
+ else
+ vars.safe_push (var);
+ }
+ if (!need_new_iterators)
+ continue;
+ if (need_new_iterators && vars.is_empty ())
+ {
+ /* No iteration variables are used in the clause - remove the
+ iterator from the clause. */
+ OMP_CLAUSE_ITERATORS (c) = NULL_TREE;
+ continue;
+ }
+
+ /* If a new iterator has been created for the current set of used
+ iterator variables, then use that as the iterator. Otherwise,
+ create a new iterator for the current iterator variable set. */
+ unsigned i;
+ for (i = 0; i < iter_vars.length (); i++)
+ {
+ if (vars.length () != iter_vars[i].length ())
+ continue;
+ bool identical_p = true;
+ for (unsigned j = 0; j < vars.length () && identical_p; j++)
+ identical_p = vars[j] == iter_vars[i][j];
+
+ if (identical_p)
+ break;
+ }
+ if (i < iter_vars.length ())
+ OMP_CLAUSE_ITERATORS (c) = new_iterators[i];
+ else
+ {
+ tree new_iters = NULL_TREE;
+ tree *new_iters_p = &new_iters;
+ tree new_vars = NULL_TREE;
+ tree *new_vars_p = &new_vars;
+ i = 0;
+ for (tree it = OMP_CLAUSE_ITERATORS (c); it && i < vars.length();
+ it = TREE_CHAIN (it))
+ {
+ tree var = TREE_VEC_ELT (it, 0);
+ if (var == vars[i])
+ {
+ *new_iters_p = copy_omp_iterator (it);
+ *new_vars_p = build_decl (OMP_CLAUSE_LOCATION (c), VAR_DECL,
+ DECL_NAME (var), TREE_TYPE (var));
+ DECL_ARTIFICIAL (*new_vars_p) = 1;
+ DECL_CONTEXT (*new_vars_p) = DECL_CONTEXT (var);
+ TREE_VEC_ELT (*new_iters_p, 0) = *new_vars_p;
+ new_iters_p = &TREE_CHAIN (*new_iters_p);
+ new_vars_p = &DECL_CHAIN (*new_vars_p);
+ i++;
+ }
+ }
+ tree new_block = make_node (BLOCK);
+ BLOCK_VARS (new_block) = new_vars;
+ TREE_VEC_ELT (new_iters, 5) = new_block;
+ new_iterators.safe_push (new_iters);
+ iter_vars.safe_push (vars.copy ());
+ OMP_CLAUSE_ITERATORS (c) = new_iters;
+ }
+
+ /* Remap clause to use the new variables. */
+ i = 0;
+ for (tree it = OMP_CLAUSE_ITERATORS (c); it; it = TREE_CHAIN (it))
+ {
+ tree old_var = vars[i++];
+ tree new_var = TREE_VEC_ELT (it, 0);
+ remap_omp_iterator_var (&OMP_CLAUSE_DECL (c), old_var, new_var);
+ remap_omp_iterator_var (&OMP_CLAUSE_SIZE (c), old_var, new_var);
+ }
+ }
+
+ for (unsigned i = 0; i < iter_vars.length (); i++)
+ iter_vars[i].release ();
+}
+
+struct iterator_loop_info_t
+{
+ tree bind;
+ tree count;
+ tree index;
+ tree body_label;
+ auto_vec<tree> clauses;
+};
+
+typedef hash_map<tree, iterator_loop_info_t> iterator_loop_info_map_t;
+
+/* 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
+ iterators are updated with information on how and where to insert code into
+ the loop body. */
+
+static void
+build_omp_iterators_loops (tree *list_p, gimple_seq *loops_seq_p)
+{
+ iterator_loop_info_map_t loops;
+
+ for (tree c = *list_p; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (!OMP_CLAUSE_HAS_ITERATORS (c))
+ continue;
+
+ bool built_p;
+ iterator_loop_info_t &loop
+ = loops.get_or_insert (OMP_CLAUSE_ITERATORS (c), &built_p);
+
+ if (!built_p)
+ {
+ loop.count = compute_omp_iterator_count (OMP_CLAUSE_ITERATORS (c),
+ loops_seq_p);
+ if (!loop.count)
+ continue;
+
+ loop.bind = NULL_TREE;
+ tree *body = build_omp_iterator_loop (OMP_CLAUSE_ITERATORS (c),
+ loops_seq_p, &loop.bind);
+
+ loop.index = create_tmp_var (sizetype);
+ SET_EXPR_LOCATION (loop.bind, OMP_CLAUSE_LOCATION (c));
+
+ /* BEFORE LOOP: */
+ /* idx = -1; */
+ /* This should be initialized to before the individual elements,
+ as idx is pre-incremented in the loop body. */
+ gimple *assign = gimple_build_assign (loop.index, size_int (-1));
+ gimple_seq_add_stmt (loops_seq_p, assign);
+
+ /* IN LOOP BODY: */
+ /* Create a label so we can find this point later. */
+ loop.body_label = create_artificial_label (OMP_CLAUSE_LOCATION (c));
+ tree tem = build1 (LABEL_EXPR, void_type_node, loop.body_label);
+ append_to_statement_list_force (tem, body);
+
+ /* idx += 2; */
+ tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+ void_type_node, loop.index,
+ size_binop (PLUS_EXPR, loop.index, size_int (2)));
+ append_to_statement_list_force (tem, body);
+ }
+
+ /* 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");
+ }
+
+ /* BEFORE LOOP: */
+ /* elems[0] = count; */
+ tree lhs = build4 (ARRAY_REF, ptr_type_node, elems, size_int (0),
+ NULL_TREE, NULL_TREE);
+ 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);
+ 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_CHAIN (new_iterator) = TREE_CHAIN (OMP_CLAUSE_ITERATORS (c));
+ OMP_CLAUSE_ITERATORS (c) = new_iterator;
+
+ loop.clauses.safe_push (c);
+ }
+
+ /* Now gimplify and add all the loops that were built. */
+ for (hash_map<tree, iterator_loop_info_t>::iterator it = loops.begin ();
+ it != loops.end (); ++it)
+ gimplify_and_add ((*it).second.bind, loops_seq_p);
+}
+
+/* Helper function for enter_omp_iterator_loop_context. */
+
+static gimple_seq *
+enter_omp_iterator_loop_context_1 (tree iterator, gimple_seq *loops_seq_p)
+{
+ /* Drill into the nested bind expressions to get to the loop body. */
+ for (gimple_stmt_iterator gsi = gsi_start (*loops_seq_p);
+ !gsi_end_p (gsi); gsi_next (&gsi))
+ {
+ gimple *stmt = gsi_stmt (gsi);
+
+ switch (gimple_code (stmt))
+ {
+ case GIMPLE_BIND:
+ {
+ gbind *bind_stmt = as_a<gbind *> (stmt);
+ gimple_push_bind_expr (bind_stmt);
+ gimple_seq *bind_body_p = gimple_bind_body_ptr (bind_stmt);
+ gimple_seq *seq =
+ enter_omp_iterator_loop_context_1 (iterator, bind_body_p);
+ if (seq)
+ return seq;
+ gimple_pop_bind_expr ();
+ }
+ break;
+ case GIMPLE_TRY:
+ {
+ gimple_seq *try_eval_p = gimple_try_eval_ptr (stmt);
+ gimple_seq *seq =
+ enter_omp_iterator_loop_context_1 (iterator, try_eval_p);
+ if (seq)
+ return seq;
+ }
+ break;
+ case GIMPLE_LABEL:
+ {
+ glabel *label_stmt = as_a<glabel *> (stmt);
+ tree label = gimple_label_label (label_stmt);
+ if (label == TREE_VEC_ELT (iterator, 6))
+ return loops_seq_p;
+ }
+ break;
+ default:
+ break;
+ }
+ }
+
+ return NULL;
+}
+
+/* Enter the Gimplification context in LOOPS_SEQ_P for the iterator loop
+ associated with OpenMP clause C. Returns the gimple_seq for the loop body
+ if C has OpenMP iterators, or ALT_SEQ_P if not. */
+
+static gimple_seq *
+enter_omp_iterator_loop_context (tree c, gimple_seq *loops_seq_p,
+ gimple_seq *alt_seq_p)
+{
+ if (!OMP_CLAUSE_HAS_ITERATORS (c))
+ return alt_seq_p;
+
+ push_gimplify_context ();
+
+ gimple_seq *seq = enter_omp_iterator_loop_context_1 (OMP_CLAUSE_ITERATORS
(c),
+ loops_seq_p);
+ gcc_assert (seq);
+ return seq;
+}
+
+/* Enter the Gimplification context in STMT for the iterator loop associated
+ with OpenMP clause C. Returns the gimple_seq for the loop body if C has
+ OpenMP iterators, or ALT_SEQ_P if not. */
+
+gimple_seq *
+enter_omp_iterator_loop_context (tree c, gomp_target *stmt,
+ gimple_seq *alt_seq_p)
+{
+ gimple_seq *loops_seq_p = gimple_omp_target_iterator_loops_ptr (stmt);
+ return enter_omp_iterator_loop_context (c, loops_seq_p, alt_seq_p);
+}
+
+/* Exit the Gimplification context for the OpenMP clause C. */
+
+void
+exit_omp_iterator_loop_context (tree c)
+{
+ if (!OMP_CLAUSE_HAS_ITERATORS (c))
+ return;
+ while (!gimplify_ctxp->bind_expr_stack.is_empty ())
+ gimple_pop_bind_expr ();
+ pop_gimplify_context (NULL);
+}
+
/* If *LIST_P contains any OpenMP depend clauses with iterators,
lower all the depend clauses by populating corresponding depend
array. Returns 0 if there are no such depend clauses, or
@@ -14924,7 +15285,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void
*data)
static void
gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
- enum tree_code code)
+ enum tree_code code,
+ gimple_seq *loops_seq_p = NULL)
{
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
tree *orig_list_p = list_p;
@@ -15295,12 +15657,14 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
: TYPE_SIZE_UNIT (TREE_TYPE (decl));
}
gimplify_omp_ctxp = ctx->outer_context;
- if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL,
+ gimple_seq *seq_p;
+ seq_p = enter_omp_iterator_loop_context (c, loops_seq_p, pre_p);
+ if (gimplify_expr (&OMP_CLAUSE_SIZE (c), seq_p, NULL,
is_gimple_val, fb_rvalue) == GS_ERROR)
{
gimplify_omp_ctxp = ctx;
remove = true;
- break;
+ goto end_adjust_omp_map_clause;
}
else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
@@ -15309,7 +15673,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
&& TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
{
OMP_CLAUSE_SIZE (c)
- = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL,
+ = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), seq_p, NULL,
false);
if ((ctx->region_type & ORT_TARGET) != 0)
omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
@@ -15350,7 +15714,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
&& (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA))
{
remove = true;
- break;
+ goto end_adjust_omp_map_clause;
}
/* If we have a DECL_VALUE_EXPR (e.g. this is a class member and/or
a variable captured in a lambda closure), look through that now
@@ -15366,7 +15730,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
decl = OMP_CLAUSE_DECL (c) = DECL_VALUE_EXPR (decl);
if (TREE_CODE (decl) == TARGET_EXPR)
{
- if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
+ if (gimplify_expr (&OMP_CLAUSE_DECL (c), seq_p, NULL,
is_gimple_lvalue, fb_lvalue) == GS_ERROR)
remove = true;
}
@@ -15453,19 +15817,19 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
/* If we have e.g. map(struct: *var), don't gimplify the
argument since omp-low.cc wants to see the decl itself. */
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT)
- break;
+ goto end_adjust_omp_map_clause;
/* We've already partly gimplified this in
gimplify_scan_omp_clauses. Don't do any more. */
if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c))
- break;
+ goto end_adjust_omp_map_clause;
gimplify_omp_ctxp = ctx->outer_context;
- if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
+ if (gimplify_expr (pd, seq_p, NULL, is_gimple_lvalue,
fb_lvalue) == GS_ERROR)
remove = true;
gimplify_omp_ctxp = ctx;
- break;
+ goto end_adjust_omp_map_clause;
}
if ((code == OMP_TARGET
@@ -15598,6 +15962,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
== GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)))
move_attach = true;
+end_adjust_omp_map_clause:
+ exit_omp_iterator_loop_context (c);
break;
case OMP_CLAUSE_TO:
@@ -18236,6 +18602,13 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq
*pre_p)
gcc_unreachable ();
}
+ gimple_seq iterator_loops_seq = NULL;
+ if (TREE_CODE (expr) == OMP_TARGET)
+ {
+ remove_unused_omp_iterator_vars (&OMP_CLAUSES (expr));
+ build_omp_iterators_loops (&OMP_CLAUSES (expr), &iterator_loops_seq);
+ }
+
bool save_in_omp_construct = in_omp_construct;
if ((ort & ORT_ACC) == 0)
in_omp_construct = false;
@@ -18279,7 +18652,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
else
gimplify_and_add (OMP_BODY (expr), &body);
gimplify_adjust_omp_clauses (pre_p, body, &OMP_CLAUSES (expr),
- TREE_CODE (expr));
+ TREE_CODE (expr), &iterator_loops_seq);
in_omp_construct = save_in_omp_construct;
switch (TREE_CODE (expr))
@@ -18322,7 +18695,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
break;
case OMP_TARGET:
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_REGION,
- OMP_CLAUSES (expr));
+ OMP_CLAUSES (expr), iterator_loops_seq);
break;
case OMP_TARGET_DATA:
/* Put use_device_{ptr,addr} clauses last, as map clauses are supposed
diff --git a/gcc/gimplify.h b/gcc/gimplify.h
index b66ceb3ce03..80c335e7c2c 100644
--- a/gcc/gimplify.h
+++ b/gcc/gimplify.h
@@ -79,6 +79,10 @@ extern enum gimplify_status gimplify_expr (tree *,
gimple_seq *, gimple_seq *,
extern tree omp_get_construct_context (void);
int omp_has_novariants (void);
+extern gimple_seq *enter_omp_iterator_loop_context (tree, gomp_target *,
+ gimple_seq * = NULL);
+extern void exit_omp_iterator_loop_context (tree);
+
extern void gimplify_type_sizes (tree, gimple_seq *);
extern void gimplify_one_sizepos (tree *, gimple_seq *);
extern gbind *gimplify_body (tree, bool);
diff --git a/gcc/gsstruct.def b/gcc/gsstruct.def
index bfe09011e55..34adc866ef2 100644
--- a/gcc/gsstruct.def
+++ b/gcc/gsstruct.def
@@ -44,6 +44,7 @@ DEFGSSTRUCT(GSS_OMP, gimple_statement_omp, false)
DEFGSSTRUCT(GSS_OMP_CRITICAL, gomp_critical, false)
DEFGSSTRUCT(GSS_OMP_FOR, gomp_for, false)
DEFGSSTRUCT(GSS_OMP_PARALLEL_LAYOUT, gimple_statement_omp_parallel_layout,
false)
+DEFGSSTRUCT(GSS_OMP_TARGET, gomp_target, false)
DEFGSSTRUCT(GSS_OMP_TASK, gomp_task, false)
DEFGSSTRUCT(GSS_OMP_SECTIONS, gomp_sections, false)
DEFGSSTRUCT(GSS_OMP_SINGLE_LAYOUT, gimple_statement_omp_single_layout, false)
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index e1036adab28..01a33c68ac7 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -12651,6 +12651,61 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
}
}
+ /* Set EXPR as the hostaddr expression that should result from the clause C.
+ LOOPS holds the intermediate loop info. Returns the tree that should be
+ passed as the hostaddr. */
+
+static tree
+lower_omp_map_iterator_expr (tree expr, tree c, gomp_target *stmt)
+{
+ if (!OMP_CLAUSE_HAS_ITERATORS (c))
+ return expr;
+
+ tree iterator = OMP_CLAUSE_ITERATORS (c);
+ tree elems = TREE_VEC_ELT (iterator, 7);
+ tree index = 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 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);
+}
+
+/* Set SIZE as the size expression that should result from the clause C.
+ LOOPS holds the intermediate loop info. Returns the tree that should be
+ passed as the clause size. */
+
+static tree
+lower_omp_map_iterator_size (tree size, tree c, gomp_target *stmt)
+{
+ if (!OMP_CLAUSE_HAS_ITERATORS (c))
+ return size;
+
+ tree iterator = OMP_CLAUSE_ITERATORS (c);
+ tree elems = TREE_VEC_ELT (iterator, 7);
+ tree index = 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 mod_expr = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+ void_type_node, lhs, size);
+ gimplify_and_add (mod_expr, loop_body_p);
+ exit_omp_iterator_loop_context (c);
+
+ return size_int (SIZE_MAX);
+}
+
/* Lower the GIMPLE_OMP_TARGET in the current statement
in GSI_P. CTX holds context information for the directive. */
@@ -12820,6 +12875,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
deep_map_cnt = extra;
}
+ if (deep_map_cnt
+ && OMP_CLAUSE_HAS_ITERATORS (c))
+ sorry ("iterators used together with deep mapping are not "
+ "supported yet");
+
if (!DECL_P (var))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
@@ -13234,6 +13294,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
*p = build_fold_indirect_ref (nd);
}
v = build_fold_addr_expr_with_type (v, ptr_type_node);
+ v = lower_omp_map_iterator_expr (v, c, stmt);
gimplify_assign (x, v, &ilist);
nc = NULL_TREE;
}
@@ -13307,13 +13368,18 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
&& TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
{
gcc_assert (offloaded);
- tree avar
- = create_tmp_var (TREE_TYPE (TREE_TYPE (x)));
- mark_addressable (avar);
- gimplify_assign (avar, build_fold_addr_expr (var), &ilist);
- talign = DECL_ALIGN_UNIT (avar);
- avar = build_fold_addr_expr (avar);
- gimplify_assign (x, avar, &ilist);
+ tree avar = build_fold_addr_expr (var);
+ if (!OMP_CLAUSE_ITERATORS (c))
+ {
+ tree tmp = create_tmp_var (TREE_TYPE (TREE_TYPE (x)));
+ mark_addressable (tmp);
+ gimplify_assign (tmp, avar, &ilist);
+ avar = tmp;
+ }
+ talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (x)));
+ avar = build_fold_addr_expr (avar);
+ avar = lower_omp_map_iterator_expr (avar, c, stmt);
+ gimplify_assign (x, avar, &ilist);
}
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
{
@@ -13392,6 +13458,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
if (s == NULL_TREE)
s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
s = fold_convert (size_type_node, s);
+ s = lower_omp_map_iterator_size (s, c, stmt);
purpose = size_int (map_idx++);
CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
if (TREE_CODE (s) != INTEGER_CST)
@@ -14324,6 +14391,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
gimple_omp_set_body (stmt, new_body);
}
+ gsi_insert_seq_before (gsi_p, gimple_omp_target_iterator_loops (stmt),
+ GSI_SAME_STMT);
+ gimple_omp_target_set_iterator_loops (stmt, NULL);
bind = gimple_build_bind (NULL, NULL,
tgt_bind ? gimple_bind_block (tgt_bind)
: NULL_TREE);
diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c
b/gcc/testsuite/c-c++-common/gomp/map-6.c
index 852839e5518..d76f9aef5aa 100644
--- a/gcc/testsuite/c-c++-common/gomp/map-6.c
+++ b/gcc/testsuite/c-c++-common/gomp/map-6.c
@@ -13,20 +13,20 @@ foo (void)
#pragma omp target map (to:a)
;
- #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type
modifier other than 'always', 'close', 'mapper' or 'present'" "" { target c++ }
} */
- ; /* { dg-error "'map' clause with map-type modifier other than 'always',
'close', 'mapper' or 'present' before 'a'" "" { target c } .-1 } */
+ #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type
modifier other than 'always', 'close', 'iterator', 'mapper' or 'present'" "" {
target c++ } } */
+ ; /* { dg-error "'map' clause with map-type modifier other than 'always',
'close', 'iterator', 'mapper' or 'present' before 'a'" "" { target c } .-1 } */
- #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with
map-type modifier other than 'always', 'close', 'mapper' or 'present'" "" {
target c++ } } */
- ; /* { dg-error "'map' clause with map-type modifier other than 'always',
'close', 'mapper' or 'present' before 'a'" "" { target c } .-1 } */
+ #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with
map-type modifier other than 'always', 'close', 'iterator', 'mapper' or
'present'" "" { target c++ } } */
+ ; /* { dg-error "'map' clause with map-type modifier other than 'always',
'close', 'iterator', 'mapper' or 'present' before 'a'" "" { target c } .-1 } */
- #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error
"'map' clause with map-type modifier other than 'always', 'close', 'mapper' or
'present'" "" { target c++ } } */
- ; /* { dg-error "'map' clause with map-type modifier other than 'always',
'close', 'mapper' or 'present' before 'a'" "" { target c } .-1 } */
+ #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error
"'map' clause with map-type modifier other than 'always', 'close', 'iterator',
'mapper' or 'present'" "" { target c++ } } */
+ ; /* { dg-error "'map' clause with map-type modifier other than 'always',
'close', 'iterator', 'mapper' or 'present' before 'a'" "" { target c } .-1 } */
- #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error
"'map' clause with map-type modifier other than 'always', 'close', 'mapper' or
'present'" "" { target c++ } } */
- ; /* { dg-error "'map' clause with map-type modifier other than 'always',
'close', 'mapper' or 'present' before 'a'" "" { target c } .-1 } */
+ #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error
"'map' clause with map-type modifier other than 'always', 'close', 'iterator',
'mapper' or 'present'" "" { target c++ } } */
+ ; /* { dg-error "'map' clause with map-type modifier other than 'always',
'close', 'iterator', 'mapper' or 'present' before 'a'" "" { target c } .-1 } */
- #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map'
clause with map-type modifier other than 'always', 'close', 'mapper' or
'present'" "" { target c++ } } */
- ; /* { dg-error "'map' clause with map-type modifier other than 'always',
'close', 'mapper' or 'present' before 'a'" "" { target c } .-1 } */
+ #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map'
clause with map-type modifier other than 'always', 'close', 'iterator',
'mapper' or 'present'" "" { target c++ } } */
+ ; /* { dg-error "'map' clause with map-type modifier other than 'always',
'close', 'iterator', 'mapper' or 'present' before 'a'" "" { target c } .-1 } */
#pragma omp target map (close a) /* { dg-error "'close' undeclared" "" {
target c } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c
b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c
new file mode 100644
index 00000000000..7d6c8dc6255
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c
@@ -0,0 +1,23 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+#define DIM1 17
+#define DIM2 39
+
+void f (int **x, int **y)
+{
+ #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2])
+ ;
+
+ #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2], y[i][:DIM2])
+ ;
+
+ #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2] + 2) /* {
dg-message "unsupported map expression" } */
+ ;
+
+ #pragma omp target map(iterator(i=0:DIM1), iterator(j=0:DIM2), to: x[i][j])
/* { dg-error "too many 'iterator' modifiers" } */
+ ;
+
+ #pragma omp target map(iterator(i=0:DIM1), to: (i % 2 == 0) ? x[i] : y[i])
/* { dg-message "unsupported map expression" } */
+ ;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c
b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c
new file mode 100644
index 00000000000..57ebb105706
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c
@@ -0,0 +1,25 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+void f (int *x, float *y, double *z)
+{
+ #pragma omp target map(iterator(i=0:10), to: x) /* { dg-warning "iterator
variable .i. not used in clause expression" } */
+ /* Add a reference to x to ensure that the 'to' clause does not get
+ dropped. */
+ x[0] = 0;
+
+ #pragma omp target map(iterator(i2=0:10, j2=0:20), from: x[i2]) /* {
dg-warning "iterator variable .j2. not used in clause expression" } */
+ ;
+
+ #pragma omp target map(iterator(i3=0:10, j3=0:20, k3=0:30), to: x[i3+j3],
y[j3+k3], z[k3+i3])
+ /* { dg-warning "iterator variable .i3. not used in clause expression" "" {
target *-*-* } .-1 } */
+ /* { dg-warning "iterator variable .j3. not used in clause expression" "" {
target *-*-* } .-2 } */
+ /* { dg-warning "iterator variable .k3. not used in clause expression" "" {
target *-*-* } .-3 } */
+ ;
+}
+
+/* { dg-final { scan-tree-dump-times "map\\\(to:x" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\\(iterator\\\(int i2=0:10:1,
loop_label=\[^\\\)\]+\\\):from:" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\\(iterator\\\(int i3=0:10:1, int
j3=0:20:1, loop_label=\[^\\\)\]+\\\):to:" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\\(iterator\\\(int j3=0:20:1, int
k3=0:30:1, loop_label=\[^\\\)\]+\\\):to:" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\\(iterator\\\(int i3=0:10:1, int
k3=0:30:1, loop_label=\[^\\\)\]+\\\):to:" 1 "gimple" } } */
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
new file mode 100644
index 00000000000..62df42ffde1
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c
@@ -0,0 +1,23 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+#define DIM1 10
+#define DIM2 20
+#define DIM3 30
+
+void f (int ***x, float ***y, double **z)
+{
+ #pragma omp target \
+ map(to: x, y) \
+ map(iterator(i=0:DIM1, j=0:DIM2), to: x[i][j][:DIM3], y[i][j][:DIM3]) \
+ map(from: z) \
+ map(iterator(i=0:DIM1), from: z[i][:DIM2])
+ ;
+}
+
+/* { 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" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-4.c
b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-4.c
new file mode 100644
index 00000000000..5dc5ad51bfb
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-4.c
@@ -0,0 +1,18 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+/* { dg-additional-options "-std=c++98" { target c++ } } */
+
+int bar (int, int);
+void baz (int, int *);
+#pragma omp declare target enter (baz)
+
+void
+foo (int x, int *p)
+{
+ #pragma omp target map (iterator (i=0:4), to: p[bar (x, i)])
+ baz (x, p);
+}
+
+/* { dg-final { scan-tree-dump "firstprivate\\\(x\\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump-times "bar \\\(x, i\\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\\(iterator\\\(int i=0:4:1,
loop_label=" 2 "gimple" } } */
diff --git a/gcc/tree-nested.cc b/gcc/tree-nested.cc
index 8d75a2f3310..5bc080caf73 100644
--- a/gcc/tree-nested.cc
+++ b/gcc/tree-nested.cc
@@ -1796,6 +1796,8 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator
*gsi, bool *handled_ops_p,
break;
case GIMPLE_OMP_TARGET:
+ walk_body (convert_nonlocal_reference_stmt,
convert_nonlocal_reference_op,
+ info, gimple_omp_target_iterator_loops_ptr (stmt));
if (!is_gimple_omp_offloaded (stmt))
{
save_suppress = info->suppress_expansion;
@@ -2517,6 +2519,9 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi,
bool *handled_ops_p,
break;
case GIMPLE_OMP_TARGET:
+ walk_body (convert_local_reference_stmt, convert_local_reference_op,
info,
+ gimple_omp_target_iterator_loops_ptr (stmt));
+
if (!is_gimple_omp_offloaded (stmt))
{
save_suppress = info->suppress_expansion;
@@ -2903,6 +2908,9 @@ convert_tramp_reference_stmt (gimple_stmt_iterator *gsi,
bool *handled_ops_p,
case GIMPLE_OMP_TASK:
do_parallel:
{
+ if (gimple_code (stmt) == GIMPLE_OMP_TARGET)
+ walk_body (convert_tramp_reference_stmt, convert_tramp_reference_op,
+ info, gimple_omp_target_iterator_loops_ptr (stmt));
tree save_local_var_chain = info->new_local_var_chain;
walk_gimple_op (stmt, convert_tramp_reference_op, wi);
info->new_local_var_chain = NULL;
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 50d08516746..663a82107f4 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -448,6 +448,15 @@ dump_omp_iterators (pretty_printer *pp, tree iter, int
spc, dump_flags_t flags)
pp_colon (pp);
dump_generic_node (pp, TREE_VEC_ELT (it, 3), spc, flags, false);
}
+ if (TREE_VEC_LENGTH (iter) > 6)
+ {
+ 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, 8), spc, flags, false);
+ }
pp_right_paren (pp);
}
@@ -1008,6 +1017,11 @@ dump_omp_clause (pretty_printer *pp, tree clause, int
spc, dump_flags_t flags)
pp_string (pp, "map(");
if (OMP_CLAUSE_MAP_READONLY (clause))
pp_string (pp, "readonly,");
+ if (OMP_CLAUSE_ITERATORS (clause))
+ {
+ dump_omp_iterators (pp, OMP_CLAUSE_ITERATORS (clause), spc, flags);
+ pp_colon (pp);
+ }
switch (OMP_CLAUSE_MAP_KIND (clause))
{
case GOMP_MAP_ALLOC:
diff --git a/gcc/tree.cc b/gcc/tree.cc
index 6a055c8c2d0..915fc9b30d4 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -323,7 +323,7 @@ unsigned const char omp_clause_num_ops[] =
1, /* OMP_CLAUSE_EXCLUSIVE */
2, /* OMP_CLAUSE_FROM */
2, /* OMP_CLAUSE_TO */
- 2, /* OMP_CLAUSE_MAP */
+ 3, /* OMP_CLAUSE_MAP */
1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */
1, /* OMP_CLAUSE_DOACROSS */
3, /* OMP_CLAUSE__MAPPER_BINDING_ */
@@ -11767,6 +11767,9 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE:
{
int len = omp_clause_num_ops[OMP_CLAUSE_CODE (t)];
+ /* Do not walk the iterator operand of OpenMP MAP clauses. */
+ if (OMP_CLAUSE_HAS_ITERATORS (t))
+ len--;
for (int i = 0; i < len; i++)
WALK_SUBTREE (OMP_CLAUSE_OPERAND (t, i));
WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (t));
diff --git a/gcc/tree.h b/gcc/tree.h
index 730b822efcc..df0b83d84e1 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1658,6 +1658,14 @@ class auto_suppress_location_wrappers
!= UNKNOWN_LOCATION)
#define OMP_CLAUSE_LOCATION(NODE) (OMP_CLAUSE_CHECK (NODE))->omp_clause.locus
+#define OMP_CLAUSE_HAS_ITERATORS(NODE) \
+ (OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_MAP \
+ && OMP_CLAUSE_ITERATORS (NODE))
+#define OMP_CLAUSE_ITERATORS(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE), \
+ OMP_CLAUSE_MAP, \
+ OMP_CLAUSE_MAP), 2)
+
/* True on OMP_FOR and other OpenMP/OpenACC looping constructs if the loop nest
is non-rectangular. */
#define OMP_FOR_NON_RECTANGULAR(NODE) \
diff --git a/libgomp/target.c b/libgomp/target.c
index cda092bd044..0ba75917143 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1003,6 +1003,105 @@ gomp_map_val (struct target_mem_desc *tgt, void
**hostaddrs, size_t i)
}
}
+static const char *
+kind_to_name (unsigned short kind)
+{
+ if (GOMP_MAP_IMPLICIT_P (kind))
+ kind &= ~GOMP_MAP_IMPLICIT;
+
+ switch (kind & 0xff)
+ {
+ case GOMP_MAP_ALLOC: return "GOMP_MAP_ALLOC";
+ case GOMP_MAP_FIRSTPRIVATE: return "GOMP_MAP_FIRSTPRIVATE";
+ case GOMP_MAP_FIRSTPRIVATE_INT: return "GOMP_MAP_FIRSTPRIVATE_INT";
+ case GOMP_MAP_TO: return "GOMP_MAP_TO";
+ case GOMP_MAP_TO_PSET: return "GOMP_MAP_TO_PSET";
+ case GOMP_MAP_FROM: return "GOMP_MAP_FROM";
+ case GOMP_MAP_TOFROM: return "GOMP_MAP_TOFROM";
+ case GOMP_MAP_POINTER: return "GOMP_MAP_POINTER";
+ case GOMP_MAP_ATTACH: return "GOMP_MAP_ATTACH";
+ case GOMP_MAP_DETACH: return "GOMP_MAP_DETACH";
+ default: return "unknown";
+ }
+}
+
+/* Map entries containing expanded iterators will be flattened and merged into
+ HOSTADDRS, SIZES and KINDS, and MAPNUM updated. Returns true if there are
+ any iterators found. ITERATOR_COUNT holds the iteration count of the
+ iterator that generates each map (0 if not generated from an iterator).
+ HOSTADDRS, SIZES, KINDS and ITERATOR_COUNT must be freed afterwards if any
+ merging occurs. */
+
+static bool
+gomp_merge_iterator_maps (size_t *mapnum, void ***hostaddrs, size_t **sizes,
+ void **kinds, size_t **iterator_count)
+{
+ bool iterator_p = false;
+ size_t map_count = 0;
+ unsigned short **skinds = (unsigned short **) kinds;
+
+ for (size_t i = 0; i < *mapnum; i++)
+ if ((*sizes)[i] == SIZE_MAX)
+ {
+ uintptr_t *iterator_array = (*hostaddrs)[i];
+ map_count += iterator_array[0];
+ iterator_p = true;
+ }
+ else
+ map_count++;
+
+ if (!iterator_p)
+ return false;
+
+ gomp_debug (1,
+ "Expanding iterator maps - number of map entries: %u -> %u\n",
+ (int) *mapnum, (int) map_count);
+ void **new_hostaddrs = (void **) gomp_malloc (map_count * sizeof (void *));
+ size_t *new_sizes = (size_t *) gomp_malloc (map_count * sizeof (size_t));
+ unsigned short *new_kinds
+ = (unsigned short *) gomp_malloc (map_count * sizeof (unsigned short));
+ size_t new_idx = 0;
+ *iterator_count = (size_t *) gomp_malloc (map_count * sizeof (size_t));
+
+ for (size_t i = 0; i < *mapnum; i++)
+ {
+ if ((*sizes)[i] == SIZE_MAX)
+ {
+ uintptr_t *iterator_array = (*hostaddrs)[i];
+ size_t count = *iterator_array++;
+ for (size_t j = 0; j < count; j++)
+ {
+ new_hostaddrs[new_idx] = (void *) *iterator_array++;
+ new_sizes[new_idx] = *iterator_array++;
+ new_kinds[new_idx] = (*skinds)[i];
+ (*iterator_count)[new_idx] = j + 1;
+ gomp_debug (1,
+ "Expanding map %u <%s>: "
+ "hostaddrs[%u] = %p, sizes[%u] = %lu\n",
+ (int) i, kind_to_name (new_kinds[new_idx]),
+ (int) new_idx, new_hostaddrs[new_idx],
+ (int) new_idx, (unsigned long) new_sizes[new_idx]);
+ new_idx++;
+ }
+ }
+ else
+ {
+ new_hostaddrs[new_idx] = (*hostaddrs)[i];
+ new_sizes[new_idx] = (*sizes)[i];
+ new_kinds[new_idx] = (*skinds)[i];
+ (*iterator_count)[new_idx] = 0;
+ new_idx++;
+ }
+ }
+
+ *mapnum = map_count;
+ *hostaddrs = new_hostaddrs;
+ *sizes = new_sizes;
+ *kinds = new_kinds;
+
+ return true;
+}
+
static inline __attribute__((always_inline)) struct target_mem_desc *
gomp_map_vars_internal (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, size_t mapnum,
@@ -1019,6 +1118,11 @@ gomp_map_vars_internal (struct gomp_device_descr
*devicep,
const int typemask = short_mapkind ? 0xff : 0x7;
struct splay_tree_s *mem_map = &devicep->mem_map;
struct splay_tree_key_s cur_node;
+ bool iterators_p = false;
+ size_t *iterator_count = NULL;
+ if (short_mapkind)
+ iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes,
+ &kinds, &iterator_count);
struct target_mem_desc *tgt
= gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
tgt->list_count = mapnum;
@@ -1896,14 +2000,17 @@ gomp_map_vars_internal (struct gomp_device_descr
*devicep,
if (pragma_kind & GOMP_MAP_VARS_TARGET)
{
+ size_t map_num = 0;
for (i = 0; i < mapnum; i++)
- {
- cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i);
- gomp_copy_host2dev (devicep, aq,
- (void *) (tgt->tgt_start + i * sizeof (void *)),
- (void *) &cur_node.tgt_offset, sizeof (void *),
- true, cbufp);
- }
+ if (!iterator_count || iterator_count[i] <= 1)
+ {
+ cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i);
+ gomp_copy_host2dev (devicep, aq,
+ (void *) (tgt->tgt_start + map_num * sizeof
(void *)),
+ (void *) &cur_node.tgt_offset, sizeof (void *),
+ true, cbufp);
+ map_num++;
+ }
}
if (cbufp)
@@ -1935,6 +2042,15 @@ gomp_map_vars_internal (struct gomp_device_descr
*devicep,
}
gomp_mutex_unlock (&devicep->lock);
+
+ if (iterators_p)
+ {
+ free (hostaddrs);
+ free (sizes);
+ free (kinds);
+ free (iterator_count);
+ }
+
return tgt;
}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c
b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c
new file mode 100644
index 00000000000..b3d87f231df
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c
@@ -0,0 +1,47 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* Test transfer of dynamically-allocated arrays to target using map
+ iterators. */
+
+#include <stdlib.h>
+
+#define DIM1 8
+#define DIM2 15
+
+int mkarray (int *x[])
+{
+ int expected = 0;
+
+ 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 expected = mkarray (x);
+
+ #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-2.c
b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c
new file mode 100644
index 00000000000..8569b55ab5b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c
@@ -0,0 +1,44 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* Test transfer of dynamically-allocated arrays from target using map
+ iterators. */
+
+#include <stdlib.h>
+
+#define DIM1 8
+#define DIM2 15
+
+void mkarray (int *x[])
+{
+ for (int i = 0; i < DIM1; i++)
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+}
+
+int main (void)
+{
+ int *x[DIM1];
+ int y, expected;
+
+ mkarray (x);
+
+ #pragma omp target enter data map(alloc: x)
+ #pragma omp target map(iterator(i=0:DIM1), from: x[i][:DIM2]) \
+ map(from: expected)
+ {
+ expected = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = (i+1) * (j+1);
+ expected += x[i][j];
+ }
+ }
+
+ 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-3.c
b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c
new file mode 100644
index 00000000000..be30fa65d80
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c
@@ -0,0 +1,56 @@
+/* { 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 and function calls in the iterator
+ expression. */
+
+#include <stdlib.h>
+
+#define DIM1 16
+#define DIM2 15
+
+int mkarrays (int *x[], int *y[])
+{
+ int expected = 0;
+
+ 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 expected = mkarrays (x, y);
+
+ #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;
+}
--
2.43.0
