This patch extends the previous patch to cover to/from clauses in
'target update'.From 99addc124535307b50fbdeb66c4f90bb0cbeb041 Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <kcye...@baylibre.com>
Date: Mon, 15 Apr 2024 13:50:22 +0100
Subject: [PATCH 3/3] openmp: Add support for iterators in to/from clauses
(C/C++)
This adds support for iterators in 'to' and 'from' clauses in the
'target update' OpenMP directive.
2024-05-24 Kwok Cheung Yeung <kcye...@baylibre.com>
gcc/c/
* c-parser.cc (c_parser_omp_clause_from_to): Parse 'iterator' modifier.
gcc/cp/
* parser.cc (cp_parser_omp_clause_from_to): Parse 'iterator' modifier.
gcc/
* gimplify.cc (gimplify_omp_map_iterators): Gimplify iterators in
to/from clauses.
(gimplify_scan_omp_clauses): Call gimplify_omp_map_iterators once to
handle clauses with iterators, then skip subsequent iterator clauses.
* omp-low.cc (scan_sharing_clauses): Skip firstprivate handling for
to/from clauses with iterators.
(lower_omp_target): Handle kinds for to/from clauses with iterators.
* tree-pretty-print.cc (dump_omp_clause): Call dump_omp_map_iterators
for to/from clauses with iterators.
gcc/testsuite/
* c-c++-common/gomp/target-update-iterator-1.c: New.
* c-c++-common/gomp/target-update-iterator-2.c: New.
* c-c++-common/gomp/target-update-iterator-3.c: New.
libgomp/
* target.c (gomp_update): Call gomp_merge_iterator_maps. Free
allocated variables.
* testsuite/libgomp.c-c++-common/target-update-iterators-1.c: New.
* testsuite/libgomp.c-c++-common/target-update-iterators-2.c: New.
* testsuite/libgomp.c-c++-common/target-update-iterators-3.c: New.
---
gcc/c/c-parser.cc | 105 ++++++++++++++--
gcc/cp/parser.cc | 116 ++++++++++++++++--
gcc/gimplify.cc | 17 ++-
gcc/omp-low.cc | 24 +++-
.../gomp/target-update-iterator-1.c | 20 +++
.../gomp/target-update-iterator-2.c | 17 +++
.../gomp/target-update-iterator-3.c | 17 +++
gcc/tree-pretty-print.cc | 20 ++-
libgomp/target.c | 12 ++
.../target-update-iterators-1.c | 65 ++++++++++
.../target-update-iterators-2.c | 57 +++++++++
.../target-update-iterators-3.c | 66 ++++++++++
12 files changed, 509 insertions(+), 27 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterator-1.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterator-2.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterator-3.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 2281148561c..6353b15d64f 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -19185,8 +19185,11 @@ c_parser_omp_clause_device_type (c_parser *parser,
tree list)
to ( variable-list )
OpenMP 5.1:
- from ( [present :] variable-list )
- to ( [present :] variable-list ) */
+ from ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
+ to ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
+
+ motion-modifier:
+ present | iterator (iterators-definition) */
static tree
c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
@@ -19197,15 +19200,88 @@ c_parser_omp_clause_from_to (c_parser *parser, enum
omp_clause_code kind,
if (!parens.require_open (parser))
return list;
+ int pos = 1, colon_pos = 0;
+ int iterator_length = 0;
+ while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)
+ {
+ if (c_parser_peek_nth_token_raw (parser, pos + 1)->type
+ == CPP_OPEN_PAREN)
+ {
+ unsigned int 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;
+ }
+ }
+ if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
+ pos += 2;
+ else
+ pos++;
+ if (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_COLON)
+ {
+ colon_pos = pos;
+ break;
+ }
+ }
+
bool present = false;
- c_token *token = c_parser_peek_token (parser);
+ tree iterators = NULL_TREE;
- if (token->type == CPP_NAME
- && strcmp (IDENTIFIER_POINTER (token->value), "present") == 0
- && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+ for (pos = 1; pos < colon_pos; pos++)
{
- present = true;
- c_parser_consume_token (parser);
+ c_token *token = c_parser_peek_token (parser);
+
+ if (token->type == CPP_COMMA)
+ {
+ c_parser_consume_token (parser);
+ continue;
+ }
+ if (token->type == CPP_NAME)
+ {
+ const char *name = IDENTIFIER_POINTER (token->value);
+ if (strcmp (name, "present") == 0)
+ {
+ if (present)
+ {
+ c_parser_error (parser, "too many %<present%> modifiers");
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ present = true;
+ c_parser_consume_token (parser);
+ }
+ else if (strcmp (name, "iterator") == 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 (kind == OMP_CLAUSE_TO)
+ c_parser_error (parser, "%<to%> clause with motion modifier "
+ "other than %<iterator%> or %<present%>");
+ else
+ c_parser_error (parser, "%<from%> clause with motion modifier "
+ "other than %<iterator%> or %<present%>");
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ }
+ }
+
+ if (colon_pos)
+ {
+ gcc_assert (pos == colon_pos);
+ gcc_assert (c_parser_next_token_is (parser, CPP_COLON));
c_parser_consume_token (parser);
}
@@ -19216,6 +19292,19 @@ c_parser_omp_clause_from_to (c_parser *parser, enum
omp_clause_code kind,
for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_MOTION_PRESENT (c) = 1;
+ if (iterators)
+ {
+ tree block = pop_scope ();
+ if (iterators == error_mark_node)
+ iterators = NULL_TREE;
+ else
+ TREE_VEC_ELT (iterators, 5) = block;
+ }
+
+ if (iterators)
+ for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+ OMP_CLAUSE_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c));
+
return nl;
}
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 6dc67851f96..f26da16ca13 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -41612,8 +41612,11 @@ cp_parser_omp_clause_doacross (cp_parser *parser, tree
list, location_t loc)
to ( variable-list )
OpenMP 5.1:
- from ( [present :] variable-list )
- to ( [present :] variable-list ) */
+ from ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
+ to ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
+
+ motion-modifier:
+ present | iterator (iterators-definition) */
static tree
cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind,
@@ -41622,15 +41625,94 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum
omp_clause_code kind,
if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
return list;
+ size_t pos = 1, colon_pos = 0;
+ int iterator_length = 0;
+ while (cp_lexer_nth_token_is (parser->lexer, pos, CPP_NAME))
+ {
+ if (cp_lexer_nth_token_is (parser->lexer, pos + 1, CPP_OPEN_PAREN))
+ {
+ unsigned int n = cp_parser_skip_balanced_tokens (parser, pos + 1);
+ if (n != pos + 1)
+ {
+ iterator_length = n - pos;
+ pos = n - 1;
+ }
+ }
+ if (cp_lexer_nth_token_is (parser->lexer, pos + 1, CPP_COMMA))
+ pos += 2;
+ else
+ pos++;
+ if (cp_lexer_nth_token_is (parser->lexer, pos, CPP_COLON))
+ {
+ colon_pos = pos;
+ break;
+ }
+ }
+
bool present = false;
- cp_token *token = cp_lexer_peek_token (parser->lexer);
+ tree iterators = NULL_TREE;
+ for (pos = 1; pos < colon_pos; pos++)
+ {
+ cp_token *token = cp_lexer_peek_token (parser->lexer);
- if (token->type == CPP_NAME
- && strcmp (IDENTIFIER_POINTER (token->u.value), "present") == 0
- && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+ if (token->type == CPP_COMMA)
+ {
+ cp_lexer_consume_token (parser->lexer);
+ continue;
+ }
+ if (token->type == CPP_NAME)
+ {
+ const char *name = IDENTIFIER_POINTER (token->u.value);
+ if (strcmp (name, "present") == 0)
+ {
+ if (present)
+ {
+ cp_parser_error (parser, "too many %<present%> modifiers");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+
/*consume_paren=*/true);
+ return list;
+ }
+ present = true;
+ cp_lexer_consume_token (parser->lexer);
+ }
+ else if (strcmp (name, "iterator") == 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 (kind == OMP_CLAUSE_TO)
+ cp_parser_error (parser, "%<to%> clause with motion modifier "
+ "other than %<iterator%> or %<present%>");
+ else
+ cp_parser_error (parser, "%<from%> clause with motion modifier "
+ "other than %<iterator%> or %<present%>");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+ }
+ }
+
+ if (colon_pos)
{
- present = true;
- cp_lexer_consume_token (parser->lexer);
+ gcc_assert (pos == colon_pos);
+ gcc_assert (cp_lexer_next_token_is (parser->lexer, CPP_COLON));
cp_lexer_consume_token (parser->lexer);
}
@@ -41639,6 +41721,19 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum
omp_clause_code kind,
for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_MOTION_PRESENT (c) = 1;
+ if (iterators)
+ {
+ tree block = poplevel (1, 1, 0);
+ if (iterators == error_mark_node)
+ iterators = NULL_TREE;
+ else
+ TREE_VEC_ELT (iterators, 5) = block;
+ }
+
+ if (iterators)
+ for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+ OMP_CLAUSE_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c));
+
return nl;
}
@@ -41684,9 +41779,8 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
&& strcmp (IDENTIFIER_POINTER (tok->u.value), "iterator") == 0
&& next_tok->type == CPP_OPEN_PAREN)
{
- size_t n = cp_parser_skip_balanced_tokens (parser, 2);
- if (cp_lexer_peek_nth_token (parser->lexer, n - 1)->type
- == CPP_CLOSE_PAREN)
+ int n = cp_parser_skip_balanced_tokens (parser, pos + 1);
+ if (n != pos + 1)
{
iterator_length = n - pos;
pos = n - 1;
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index cadb2a3d96d..3cf88cf82b5 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9304,7 +9304,9 @@ gimplify_omp_map_iterators (tree *list_p, gimple_seq
*pre_p)
while (tree c = *list_p)
{
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_TO
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FROM)
{
list_p = &OMP_CLAUSE_CHAIN (c);
continue;
@@ -12937,6 +12939,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
case OMP_CLAUSE_FROM:
case OMP_CLAUSE__CACHE_:
decl = OMP_CLAUSE_DECL (c);
+
+ if (OMP_ITERATOR_DECL_P (decl))
+ {
+ if (!handled_map_iterators)
+ {
+ gimplify_omp_map_iterators (list_p, pre_p);
+ handled_map_iterators = true;
+ continue;
+ }
+ /* Skip declarations with iterators. */
+ break;
+ }
+
if (error_operand_p (decl))
{
remove = true;
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index d0f0c5d884a..9e94ad329cd 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1519,8 +1519,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
{
/* FIXME: Is this the right way to handle these? */
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
break;
tree field
= build_decl (OMP_CLAUSE_LOCATION (c),
@@ -13061,7 +13063,23 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
purpose = size_int (map_idx++);
CONSTRUCTOR_APPEND_ELT (vsize, purpose, size_int (SIZE_MAX));
- unsigned HOST_WIDE_INT tkind = OMP_CLAUSE_MAP_KIND (c);
+ unsigned HOST_WIDE_INT tkind = 0;
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_TO:
+ tkind = (OMP_CLAUSE_MOTION_PRESENT (c)
+ ? GOMP_MAP_ALWAYS_PRESENT_TO : GOMP_MAP_TO);
+ break;
+ case OMP_CLAUSE_FROM:
+ tkind = (OMP_CLAUSE_MOTION_PRESENT (c)
+ ? GOMP_MAP_ALWAYS_PRESENT_FROM : GOMP_MAP_FROM);
+ break;
+ case OMP_CLAUSE_MAP:
+ tkind = OMP_CLAUSE_MAP_KIND (c);
+ break;
+ default:
+ gcc_unreachable ();
+ }
gcc_checking_assert (tkind
< (HOST_WIDE_INT_C (1U) << talign_shift));
gcc_checking_assert (tkind
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterator-1.c
b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-1.c
new file mode 100644
index 00000000000..3a64f511da4
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-1.c
@@ -0,0 +1,20 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+#define DIM1 17
+#define DIM2 39
+
+void f (int **x, float **y)
+{
+ #pragma omp target update to (iterator(i=0:DIM1): x[i][:DIM2])
+
+ #pragma omp target update to (iterator(i=0:DIM1): x[i][:DIM2], y[i][:DIM2])
+
+ #pragma omp target update to (iterator(i=0:DIM1), present: x[i][:DIM2])
+
+ #pragma omp target update to (iterator(i=0:DIM1), iterator(j=0:DIM2):
x[i][j]) /* { dg-error "too many 'iterator' modifiers" } */
+ /* { dg-error ".#pragma omp target update. must contain at least one .from.
or .to. clauses" "" { target *-*-* } .-1 } */
+
+ #pragma omp target update to (iterator(i=0:DIM1), something: x[i][j]) /* {
dg-error ".to. clause with motion modifier other than .iterator. or .present.
before .something." } */
+ /* { dg-error ".#pragma omp target update. must contain at least one .from.
or .to. clauses" "" { target *-*-* } .-1 } */
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterator-2.c
b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-2.c
new file mode 100644
index 00000000000..3789a559b6f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-2.c
@@ -0,0 +1,17 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+void f (int *x, float *y, double *z)
+{
+ #pragma omp target update to(iterator(i=0:10): x) /* { dg-error "iterator
variable .i. not used in clause expression" }*/
+ ;
+
+ #pragma omp target update from(iterator(i=0:10, j=0:20): x[i]) /* { dg-error
"iterator variable .j. not used in clause expression" }*/
+ ;
+
+ #pragma omp target update to(iterator(i=0:10, j=0:20, k=0:30): x[i], y[j],
z[k])
+ /* { dg-error "iterator variable .i. not used in clause expression" "" {
target *-*-* } .-1 } */
+ /* { dg-error "iterator variable .j. not used in clause expression" "" {
target *-*-* } .-2 } */
+ /* { dg-error "iterator variable .k. not used in clause expression" "" {
target *-*-* } .-3 } */
+ ;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterator-3.c
b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-3.c
new file mode 100644
index 00000000000..d8672b3a242
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-3.c
@@ -0,0 +1,17 @@
+/* { 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 update to (iterator(i=0:DIM1, j=0:DIM2): x[i][j][:DIM3],
y[i][j][:DIM3])
+ #pragma omp target update from (iterator(i=0:DIM1): z[i][:DIM2])
+}
+
+/* { dg-final { scan-tree-dump-times "if \\(i <= 9\\) goto <D\.\[0-9\]+>; else
goto <D\.\[0-9\]+>;" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto <D\.\[0-9\]+>;
else goto <D\.\[0-9\]+>;" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "to\\(iterator\\(int i=0:10:1, int
j=0:20:1\\):iterator_array=D\.\[0-9\]+" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "from\\(iterator\\(int
i=0:10:1\\):iterator_array=D\.\[0-9\]+" 1 "gimple" } } */
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index feb9c93a809..a0fe270ab1b 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1099,16 +1099,28 @@ dump_omp_clause (pretty_printer *pp, tree clause, int
spc, dump_flags_t flags)
pp_string (pp, "from(");
if (OMP_CLAUSE_MOTION_PRESENT (clause))
pp_string (pp, "present:");
- dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
- spc, flags, false);
+ decl = OMP_CLAUSE_DECL (clause);
+ if (OMP_ITERATOR_DECL_P (decl))
+ {
+ dump_omp_map_iterators (pp, TREE_PURPOSE (decl), spc, flags);
+ pp_colon (pp);
+ decl = TREE_VALUE (decl);
+ }
+ dump_generic_node (pp, decl, spc, flags, false);
goto print_clause_size;
case OMP_CLAUSE_TO:
pp_string (pp, "to(");
if (OMP_CLAUSE_MOTION_PRESENT (clause))
pp_string (pp, "present:");
- dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
- spc, flags, false);
+ decl = OMP_CLAUSE_DECL (clause);
+ if (OMP_ITERATOR_DECL_P (decl))
+ {
+ dump_omp_map_iterators (pp, TREE_PURPOSE (decl), spc, flags);
+ pp_colon (pp);
+ decl = TREE_VALUE (decl);
+ }
+ dump_generic_node (pp, decl, spc, flags, false);
goto print_clause_size;
case OMP_CLAUSE__CACHE_:
diff --git a/libgomp/target.c b/libgomp/target.c
index 07ba840b495..71a6428af49 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2221,6 +2221,7 @@ gomp_update (struct gomp_device_descr *devicep, size_t
mapnum, void **hostaddrs,
size_t i;
struct splay_tree_key_s cur_node;
const int typemask = short_mapkind ? 0xff : 0x7;
+ bool iterators_p = false;
if (!devicep)
return;
@@ -2228,6 +2229,10 @@ gomp_update (struct gomp_device_descr *devicep, size_t
mapnum, void **hostaddrs,
if (mapnum == 0)
return;
+ if (short_mapkind)
+ iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes,
+ &kinds);
+
gomp_mutex_lock (&devicep->lock);
if (devicep->state == GOMP_DEVICE_FINALIZED)
{
@@ -2321,6 +2326,13 @@ gomp_update (struct gomp_device_descr *devicep, size_t
mapnum, void **hostaddrs,
}
}
gomp_mutex_unlock (&devicep->lock);
+
+ if (iterators_p)
+ {
+ free (hostaddrs);
+ free (sizes);
+ free (kinds);
+ }
}
static struct gomp_offload_icv_list *
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c
b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c
new file mode 100644
index 00000000000..5a4cad5c219
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c
@@ -0,0 +1,65 @@
+/* { dg-do run } */
+
+/* Test target enter data and target update to the 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 sum;
+ int expected = mkarray (x);
+
+ #pragma omp target enter data map(to: x[:DIM1])
+ #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2])
+ #pragma omp target map(from: sum)
+ {
+ sum = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j];
+ }
+
+ if (sum != expected)
+ return 1;
+
+ expected = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] *= rand ();
+ expected += x[i][j];
+ }
+
+ #pragma omp target update to(iterator(i=0:DIM1): x[i][:DIM2])
+
+ #pragma omp target map(from: sum)
+ {
+ sum = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j];
+ }
+
+ return sum != expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c
b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c
new file mode 100644
index 00000000000..949cc266d84
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c
@@ -0,0 +1,57 @@
+/* { dg-do run } */
+
+/* Test target enter data and target update from the 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));
+ for (int j = 0; j < DIM2; j++)
+ x[i][j] = 0;
+ }
+}
+
+int main (void)
+{
+ int *x[DIM1];
+ int sum, expected;
+
+ mkarray (x);
+
+ #pragma omp target enter data map(alloc: x[:DIM1])
+ #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2])
+ #pragma omp target 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 + 2);
+ expected += x[i][j];
+ }
+ }
+
+ /* Host copy of x should remain unchanged. */
+ sum = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j];
+ if (sum != 0)
+ return 1;
+
+ #pragma omp target update from(iterator(i=0:DIM1): x[i][:DIM2])
+
+ /* Host copy should now be updated. */
+ sum = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j];
+ return sum - expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c
b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c
new file mode 100644
index 00000000000..852635e50f6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c
@@ -0,0 +1,66 @@
+/* { dg-do run } */
+
+/* Test target enter data and target update to the target using map
+ iterators with a function. */
+
+#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));
+ for (int j = 0; j < DIM2; j++)
+ x[i][j] = rand ();
+ }
+}
+
+int f (int i)
+{
+ return i * 2;
+}
+
+int main (void)
+{
+ int *x[DIM1], x_new[DIM1][DIM2];
+ int sum, expected;
+
+ mkarray (x);
+
+ #pragma omp target enter data map(alloc: x[:DIM1])
+ #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2])
+
+ /* Update x on host. */
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ {
+ x_new[i][j] = x[i][j];
+ x[i][j] = (i + 1) * (j + 2);
+ }
+
+ /* Update a subset of x on target. */
+ #pragma omp target update to(iterator(i=0:DIM1/2): x[f (i)][:DIM2])
+
+ #pragma omp target map(from: sum)
+ {
+ sum = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j];
+ }
+
+ /* Calculate expected value on host. */
+ for (int i = 0; i < DIM1/2; i++)
+ for (int j = 0; j < DIM2; j++)
+ x_new[f (i)][j] = x[f (i)][j];
+
+ expected = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ expected += x_new[i][j];
+
+ return sum - expected;
+}
--
2.34.1
