Hi Tobias,
Thanks for the review.
On 17/07/2025 10:09 am, Tobias Burnus wrote:
Regarding gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c:
target-map-iterators-2.c:6:25: warning: iterator variable ‘i’ not used
in clause expression
and
+ warning_at (OMP_CLAUSE_LOCATION (c), 0,
+ "iterator variable %qE not used in clause "
+ "expression", DECL_NAME (var));
Can you replace '0' by some OPT_… such as OPT_Wopenmp?
Done.
Additionally, while I think it is okay to have a warning for:
#pragma omp target map(iterator(i2=0:10, j2=0:20), from: x[i2])
/* { dg-warning "iterator variable .j2. not used in clause
expression" } */
I find the following warning odd:
#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 *-*-* } .-2 } */
/* { dg-warning "iterator variable .j3. not used in clause
expression" "" { target *-*-* } .-3 } */
/* { dg-warning "iterator variable .k3. not used in clause
expression" "" { target *-*-* } .-4 } */
as all variables appears in expressions!
In the original dump, there is:
map(iterator(int i3=0:10:1, int j3=0:20:1, int k3=0:30:1):
to:*(z + (sizetype) ((long unsigned int) (k3 + i3) * 8)) [len: 8])
I understand that the iterator can be removed as x/y/z are unused, but this
does not make the warning more sensible, either.
This is due to there being multiple maps in a single clause at the
source level. The FE allocates a separate clause for each map, but for
some reason the C FE always sets the clause location to after the 'map'
for all clauses. I've updated the clause location to point to the map
expression (which is what the C++ and Fortran FEs do) which hopefully
should be clearer.
target-map-iterators-2.c:14:87: warning: iterator variable ‘j3’ not used
in clause expression [-Wopenmp]
14 | #pragma omp target map(iterator(i3=0:10, j3=0:20, k3=0:30),
to: x[i3+j3], y[j3+k3], z[k3+i3])
|
^
target-map-iterators-2.c:14:77: warning: iterator variable ‘i3’ not used
in clause expression [-Wopenmp]
14 | #pragma omp target map(iterator(i3=0:10, j3=0:20, k3=0:30),
to: x[i3+j3], y[j3+k3], z[k3+i3])
|
^
target-map-iterators-2.c:14:67: warning: iterator variable ‘k3’ not used
in clause expression [-Wopenmp]
14 | #pragma omp target map(iterator(i3=0:10, j3=0:20, k3=0:30),
to: x[i3+j3], y[j3+k3], z[k3+i3])
|
^
I also experimented with actually printing out the expression in the
warning, but the FEs convert these to a pointer expression long before
it reaches the ME, so it would probably be even more confusing!
* * *
BTW: I wonder whether we should special case:
(a) begin:end with begin == end → single iteration
(b) begin:end with begin < end (zero-sized empty loop)
And we may want to warn for step == 0.
Currently the algorithm for adding code to the iterator loop after it
has been built relies on the presence of the loop, so removing the loop
for 0-1 iterations would complicate matters. Hopefully the dead-code
elimination can remove any unrequired loops, but if not it could be
added as an optimisation later.
step == 0 currently results in an ICE, but that is in generic iterator
code and should probably be dealt with separately to this patch.
In any case, it would be good to have a testcase that checks for
empty loops, i.e. begin > end (with step < 0) and begin > end (with step
< 0)
And, I guess, it wouldn't harm to have a check for step < 0 with begin >
end
(finite loop size but with negative values).
I've added tests in gomp/target-map-iterators-2.c. I've also added a
warning if the number of iterations can be statically determined to be zero.
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
I think we want to have somewhere a comment regarding how iterators are
passed. I think the main one should be in libgomp/target.c but possibly
some hint about the format makes sense. For instance, as part of the
comment to (and/or) lower_omp_map_iterator_{expr,size}.
I've added comments to lower_omp_map_iterator_* regarding what they return.
+ 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);
Twice: Line starts with space + tab.
Fixed.
case GIMPLE_OMP_TARGET:
if (!is_gimple_omp_offloaded (stmt))
{
*handled_ops_p = false;
return NULL_TREE;
}
/* FALLTHRU */
case GIMPLE_OMP_PARALLEL:
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));
I think it would be cleaner to move the walk_body before FALLTHRU.
Done.
--- 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--;
This looks fragile, i.e. it assumes that we never add a new operator to
OMP_CLAUSE_MAP; I wonder whether should add somewhere a note about this?
I don't want to clutter omp_clause_num_ops but it seems to be the safest
place?
I've added a small comment to omp_clause_num_ops. Though I don't think
you would have to change anything in walk_tree_1 as long as you always
keep the iterator as the last operand?
--- 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";
+ }
+}
I think that's a quite useful function; I also have locally a similar
function
to debug mapping issues. However, I think you should really cover all
kinds,
e.g. I fail to see why 'map(iterator(), always, to:' should print
'unknown'.
If I grep correctly, I see the following GOMP_MAP items being used in
target;
I have not checked whether all make sense, though - at least implicit
can be
removed for the current code.
Probably, we should then also add the short flag to the function. In the
caller we can pass the fixed value ('true'), but that permits to later
reuse
the function simply (locally or also in a committed version) for other map
kinds.
GOMP_MAP_ALLOC
GOMP_MAP_ALWAYS_FROM
GOMP_MAP_ALWAYS_POINTER
GOMP_MAP_ALWAYS_PRESENT_FROM
GOMP_MAP_ALWAYS_PRESENT_TO
GOMP_MAP_ALWAYS_PRESENT_TOFROM
GOMP_MAP_ALWAYS_TO
GOMP_MAP_ALWAYS_TOFROM
GOMP_MAP_ATTACH
GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION
GOMP_MAP_DELETE
GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
GOMP_MAP_DETACH
GOMP_MAP_FIRSTPRIVATE
GOMP_MAP_FIRSTPRIVATE_INT
GOMP_MAP_FORCE_ALLOC
GOMP_MAP_FORCE_DEVICEPTR
GOMP_MAP_FORCE_FROM
GOMP_MAP_FORCE_PRESENT
GOMP_MAP_FORCE_TO
GOMP_MAP_FORCE_TOFROM
GOMP_MAP_FROM
GOMP_MAP_IF_PRESENT
GOMP_MAP_IMPLICIT
GOMP_MAP_POINTER
GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION
GOMP_MAP_RELEASE
GOMP_MAP_STRUCT
GOMP_MAP_STRUCT_UNORD
GOMP_MAP_TO
GOMP_MAP_TOFROM
GOMP_MAP_TO_PSET
GOMP_MAP_USE_DEVICE_PTR
GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT
GOMP_MAP_VARS_DATA
GOMP_MAP_VARS_ENTER_DATA
GOMP_MAP_VARS_OPENACC
GOMP_MAP_VARS_TARGET
GOMP_MAP_ZERO_LEN_ARRAY_SECTION
I've now extended kind_to_name to handle all map types other than the
internal ones (which shouldn't reach libgomp). I have also added a
parameter to indicate whether the kind is short.
+/* 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).
(To me it reads better with an 'and': '… each map (and 0 if …)'.)
Done.
+ HOSTADDRS, SIZES, KINDS and ITERATOR_COUNT must be freed
afterwards if any
+ merging occurs. */
I think we should somewhere in target.c document the format of
iterators, either here
or for the call, i.e. SIZE == MAX_SIZE to indicate the iterator, kind
contains the
size and hostpointer points to an array of triplets (addr,kind,size) for
each
iterated item.
That's hidden in the code, but it helps to have it spelled out to make
debugging
of general mapping issues easier.
I've added an explanation regarding how map iterators are passed to
libgomp in the comments to gomp_merge_iterator_maps.
@@ -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);
I think it makes sense to append ' /* OpenMP */' to the 'if
(short_mapkind)' line.
Done.
@@ -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++;
+ }
}
First, I think it is more readable with '== 1' instead of '<= 1' as == 0
cannot
occur. Second, I think that requires a comment why > 1 is ignored here.
iterator_count[i] can be 0 if there are maps with iterators present, but
the i'th map happens to not be using one (which is why the count starts
at 1 rather than 0 for iterator maps). We don't really use this at the
moment, but there is no harm in providing an additional way to
differentiate between 'normal' maps and iterator-generated maps.
I have added a comment explaining why iterator maps after the first must
be skipped.
* * *
Thanks for the patch,
Tobias
Thanks
Kwok
From 409da06c185354bc50ad13b67af52359be8ba7e4 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] 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_variable_list): Use location of the
map expression as the clause location.
(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 | 56 ++-
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 | 401 +++++++++++++++++-
gcc/gimplify.h | 4 +
gcc/gsstruct.def | 1 +
gcc/omp-low.cc | 82 +++-
gcc/testsuite/c-c++-common/gomp/map-6.c | 20 +-
.../gomp/target-map-iterators-1.c | 23 +
.../gomp/target-map-iterators-2.c | 41 ++
.../gomp/target-map-iterators-3.c | 23 +
.../gomp/target-map-iterators-4.c | 18 +
gcc/tree-nested.cc | 7 +
gcc/tree-pretty-print.cc | 14 +
gcc/tree.cc | 5 +-
gcc/tree.h | 8 +
libgomp/target.c | 185 +++++++-
.../target-map-iterators-1.c | 47 ++
.../target-map-iterators-2.c | 44 ++
.../target-map-iterators-3.c | 56 +++
25 files changed, 1132 insertions(+), 56 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 5119841a589..7858f351dfd 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -16659,7 +16659,7 @@ c_parser_omp_variable_list (c_parser *parser,
|| CONVERT_EXPR_P (decl))
decl = TREE_OPERAND (decl, 0);
- tree u = build_omp_clause (clause_loc, kind);
+ tree u = build_omp_clause (loc, kind);
OMP_CLAUSE_DECL (u) = decl;
OMP_CLAUSE_CHAIN (u) = list;
list = u;
@@ -20072,7 +20072,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)
@@ -20087,15 +20087,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)
@@ -20117,6 +20137,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);
@@ -20150,6 +20171,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);
@@ -20223,8 +20255,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;
}
@@ -20273,9 +20305,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 f161bd9d0e7..5988c6ff5a6 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -16195,7 +16195,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))
{
@@ -16950,6 +16957,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:
@@ -17681,6 +17694,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 0d9ed2ea82b..08691e99e70 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -42598,16 +42598,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)
@@ -42620,6 +42638,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);
@@ -42658,6 +42677,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);
@@ -42746,7 +42780,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,
@@ -42810,9 +42844,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 86b09049677..7926434cb63 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -7743,7 +7743,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))
{
@@ -8992,6 +8999,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:
@@ -9898,6 +9911,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 5c970cee7e6..da32651ea01 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 910314bd54c..1bdda6728a2 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9781,6 +9781,371 @@ 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_TO
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM))
+ warning_at (OMP_CLAUSE_LOCATION (c), OPT_Wopenmp,
+ "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;
+ if (TREE_CONSTANT (loop.count) && integer_zerop (loop.count))
+ warning_at (OMP_CLAUSE_LOCATION (c), OPT_Wopenmp,
+ "iteration count is zero");
+
+ 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
@@ -14925,7 +15290,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;
@@ -15296,12 +15662,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)
@@ -15310,7 +15678,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),
@@ -15351,7 +15719,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
@@ -15367,7 +15735,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;
}
@@ -15454,19 +15822,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
@@ -15599,6 +15967,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:
@@ -18237,6 +18607,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;
@@ -18280,7 +18657,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))
@@ -18323,7 +18700,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..9d80a3573dc 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -12651,6 +12651,63 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
}
}
+ /* Set EXPR as the hostaddr expression that should result from the clause C
+ in the target statement STMT. Returns the tree that should be
+ passed as the hostaddr (a pointer to the array containing the expanded
+ hostaddrs and sizes of the clause). */
+
+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
+ in the target statement STMT. Returns the tree that should be
+ passed as the clause size (a size_int with the value SIZE_MAX, indicating
+ that the clause uses an iterator). */
+
+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 +12877,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 +13296,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,12 +13370,17 @@ 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);
+ 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 +13460,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 +14393,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..e151baf2d08
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c
@@ -0,0 +1,41 @@
+/* { 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 } */
+ ;
+
+ /* Test iterator with zero iterations. */
+ #pragma omp target map(iterator(i4=0:0), to: x[i4]) /* { dg-warning
"iteration count is zero" } */
+ ;
+
+ /* Test iterator where the beginning is greater than the end. */
+ #pragma omp target map(iterator(i5=10:0), to: x[i5]) /* { dg-warning
"iteration count is zero" } */
+ ;
+
+ /* Test iterator where the beginning is greater than the end, but with a
+ negative step. */
+ #pragma omp target map(iterator(i6=10:0:-1), to: x[i6])
+ ;
+}
+
+/* { 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" } } */
+/* { dg-final { scan-tree-dump-times "map\\\(iterator\\\(int i4=0:0:1,
loop_label=\[^\\\)\]+\\\):to:" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\\(iterator\\\(int i5=10:0:1,
loop_label=\[^\\\)\]+\\\):to:" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\\(iterator\\\(int i6=10:0:-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..813334b5c92 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;
@@ -2898,6 +2903,8 @@ convert_tramp_reference_stmt (gimple_stmt_iterator *gsi,
bool *handled_ops_p,
*handled_ops_p = false;
return NULL_TREE;
}
+ walk_body (convert_tramp_reference_stmt, convert_tramp_reference_op,
+ info, gimple_omp_target_iterator_loops_ptr (stmt));
/* FALLTHRU */
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
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 0f02924763f..203cc1ccd7a 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -325,7 +325,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 (update walk_tree_1 if this is changed) */
1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */
1, /* OMP_CLAUSE_DOACROSS */
3, /* OMP_CLAUSE__MAPPER_BINDING_ */
@@ -11769,6 +11769,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 c0e434ba897..438065738db 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1659,6 +1659,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..43976f0195b 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1003,6 +1003,155 @@ gomp_map_val (struct target_mem_desc *tgt, void
**hostaddrs, size_t i)
}
}
+static const char *
+kind_to_name (unsigned short kind, bool short_mapkind)
+{
+ if (short_mapkind && GOMP_MAP_IMPLICIT_P (kind))
+ kind &= ~GOMP_MAP_IMPLICIT;
+
+ switch (kind & (short_mapkind ? 0xff : 0x7))
+ {
+ case GOMP_MAP_ALLOC: return "GOMP_MAP_ALLOC";
+ case GOMP_MAP_TO: return "GOMP_MAP_TO";
+ 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_TO_PSET: return "GOMP_MAP_TO_PSET";
+ case GOMP_MAP_FORCE_PRESENT: return "GOMP_MAP_FORCE_PRESENT";
+ case GOMP_MAP_DELETE: return "GOMP_MAP_DELETE";
+ case GOMP_MAP_FORCE_DEVICEPTR: return "GOMP_MAP_FORCE_DEVICEPTR";
+ case GOMP_MAP_DEVICE_RESIDENT: return "GOMP_MAP_DEVICE_RESIDENT";
+ case GOMP_MAP_LINK: return "GOMP_MAP_LINK";
+ case GOMP_MAP_IF_PRESENT: return "GOMP_MAP_IF_PRESENT";
+ case GOMP_MAP_FIRSTPRIVATE: return "GOMP_MAP_FIRSTPRIVATE";
+ case GOMP_MAP_FIRSTPRIVATE_INT: return "GOMP_MAP_FIRSTPRIVATE_INT";
+ case GOMP_MAP_USE_DEVICE_PTR: return "GOMP_MAP_USE_DEVICE_PTR";
+ case GOMP_MAP_ZERO_LEN_ARRAY_SECTION: return
"GOMP_MAP_ZERO_LEN_ARRAY_SECTION";
+ case GOMP_MAP_FORCE_ALLOC: return "GOMP_MAP_FORCE_ALLOC";
+ case GOMP_MAP_FORCE_TO: return "GOMP_MAP_FORCE_TO";
+ case GOMP_MAP_FORCE_FROM: return "GOMP_MAP_FORCE_FROM";
+ case GOMP_MAP_FORCE_TOFROM: return "GOMP_MAP_FORCE_TOFROM";
+ case GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT:
+ return "GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT";
+ case GOMP_MAP_ALWAYS_TO: return "GOMP_MAP_ALWAYS_TO";
+ case GOMP_MAP_ALWAYS_FROM: return "GOMP_MAP_ALWAYS_FROM";
+ case GOMP_MAP_ALWAYS_TOFROM: return "GOMP_MAP_ALWAYS_TOFROM";
+ case GOMP_MAP_ALWAYS_PRESENT_TO: return "GOMP_MAP_ALWAYS_PRESENT_TO";
+ case GOMP_MAP_ALWAYS_PRESENT_FROM: return "GOMP_MAP_ALWAYS_PRESENT_FROM";
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM: return
"GOMP_MAP_ALWAYS_PRESENT_TOFROM";
+ case GOMP_MAP_STRUCT: return "GOMP_MAP_STRUCT";
+ case GOMP_MAP_STRUCT_UNORD: return "GOMP_MAP_STRUCT_UNORD";
+ case GOMP_MAP_ALWAYS_POINTER: return "GOMP_MAP_ALWAYS_POINTER";
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+ return "GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION";
+ case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION:
+ return "GOMP_MAP_DELETE_ZERO_LENGTH_ARRAY_SECTION";
+ case GOMP_MAP_RELEASE: return "GOMP_MAP_RELEASE";
+ case GOMP_MAP_ATTACH: return "GOMP_MAP_ATTACH";
+ case GOMP_MAP_DETACH: return "GOMP_MAP_DETACH";
+ case GOMP_MAP_FORCE_DETACH: return "GOMP_MAP_FORCE_DETACH";
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+ return "GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION";
+ default: return "unknown";
+ }
+}
+
+/* When GCC encounters a clause with an iterator, e.g.:
+
+ #pragma omp target map (iterator(i=0:4), to: x[i])
+
+ it generates an array containing the number of iterations and the
+ address and size of each iteration. e.g.:
+
+ void *omp_iter_data[] = {
+ (void *) 4, // Number of iterations
+ &x[0], (void *) sizeof(x[0]),
+ &x[1], (void *) sizeof(x[1]),
+ &x[2], (void *) sizeof(x[2]),
+ &x[3], (void *) sizeof(x[3])
+ };
+
+ When the construct is lowered, &omp_iter_data is used as the host address
+ for the map (instead of &x[i]), and the size is set to SIZE_MAX to mark
+ the map as an iterator map.
+
+ 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 (and 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], true),
+ (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 +1168,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) /* OpenMP */
+ 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 +2050,22 @@ gomp_map_vars_internal (struct gomp_device_descr
*devicep,
if (pragma_kind & GOMP_MAP_VARS_TARGET)
{
+ /* The target variables table is constructed with maps using iterators
+ unexpanded. Now that the iterator maps are expanded, we will need to
+ skip all expanded maps after the initial entry, otherwise subsequent
+ maps will be out-of-sync with their corresponding entry in the
+ target variables table. */
+ 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 +2097,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
