https://gcc.gnu.org/g:3e3966d78461a4cfadfde079812e3d98335095b0
commit 3e3966d78461a4cfadfde079812e3d98335095b0
Author: Kwok Cheung Yeung <kcye...@baylibre.com>
Date: Wed Nov 27 21:51:34 2024 +0000
openmp: Add support for iterators in 'target update' clauses (C/C++)
This adds support for iterators in 'to' and 'from' clauses in the
'target update' OpenMP directive.
gcc/c/
* c-parser.cc (c_parser_omp_clause_from_to): Parse 'iterator'
modifier.
* c-typeck.cc (c_finish_omp_clauses): Finish iterators for to/from
clauses.
gcc/cp/
* parser.cc (cp_parser_omp_clause_from_to): Parse 'iterator'
modifier.
* semantics.cc (finish_omp_clauses): Finish iterators for to/from
clauses.
gcc/
* gimplify.cc (gimplify_scan_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): Add argument for iterator loops sequence
in call to gimplify_scan_omp_clauses.
(gimplify_omp_target_update): Call remove_unused_omp_iterator_vars
and
build_omp_iterators_loops. Add loop sequence as argument when
calling
gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses and building
the Gimple statement.
* tree-pretty-print.cc (dump_omp_clause): Call dump_omp_iterators
for to/from clauses with iterators.
* tree.cc (omp_clause_num_ops): Add extra operand for
OMP_CLAUSE_FROM
and OMP_CLAUSE_TO.
* tree.h (OMP_CLAUSE_HAS_ITERATORS): Add check for OMP_CLAUSE_TO and
OMP_CLAUSE_FROM.
(OMP_CLAUSE_ITERATORS): Likewise.
gcc/testsuite/
* c-c++-common/gomp/target-update-iterators-1.c: New.
* c-c++-common/gomp/target-update-iterators-2.c: New.
* c-c++-common/gomp/target-update-iterators-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.
Diff:
---
gcc/ChangeLog.omp | 19 ++++++
gcc/c/ChangeLog.omp | 6 ++
gcc/c/c-parser.cc | 55 ++++++++++++++----
gcc/c/c-typeck.cc | 5 +-
gcc/cp/ChangeLog.omp | 6 ++
gcc/cp/parser.cc | 56 ++++++++++++++++--
gcc/cp/semantics.cc | 5 +-
gcc/gimplify.cc | 37 +++++++-----
gcc/testsuite/ChangeLog.omp | 6 ++
.../c-c++-common/gomp/target-update-iterators-1.c | 20 +++++++
.../c-c++-common/gomp/target-update-iterators-2.c | 23 ++++++++
.../c-c++-common/gomp/target-update-iterators-3.c | 17 ++++++
gcc/tree-pretty-print.cc | 10 ++++
gcc/tree.cc | 4 +-
gcc/tree.h | 6 +-
libgomp/ChangeLog.omp | 8 +++
libgomp/target.c | 14 +++++
.../target-update-iterators-1.c | 65 +++++++++++++++++++++
.../target-update-iterators-2.c | 58 +++++++++++++++++++
.../target-update-iterators-3.c | 67 ++++++++++++++++++++++
20 files changed, 449 insertions(+), 38 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 0d31c77933e0..2ba1ef83f204 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,22 @@
+2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com>
+
+ * gimplify.cc (gimplify_scan_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): Add argument for iterator loops sequence
+ in call to gimplify_scan_omp_clauses.
+ (gimplify_omp_target_update): Call remove_unused_omp_iterator_vars and
+ build_omp_iterators_loops. Add loop sequence as argument when calling
+ gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses and building
+ the Gimple statement.
+ * tree-pretty-print.cc (dump_omp_clause): Call dump_omp_iterators
+ for to/from clauses with iterators.
+ * tree.cc (omp_clause_num_ops): Add extra operand for OMP_CLAUSE_FROM
+ and OMP_CLAUSE_TO.
+ * tree.h (OMP_CLAUSE_HAS_ITERATORS): Add check for OMP_CLAUSE_TO and
+ OMP_CLAUSE_FROM.
+ (OMP_CLAUSE_ITERATORS): Likewise.
+
2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com>
* gimple-pretty-print.cc (dump_gimple_omp_target): Print expanded
diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp
index cc3a28c836db..af995a398571 100644
--- a/gcc/c/ChangeLog.omp
+++ b/gcc/c/ChangeLog.omp
@@ -1,3 +1,9 @@
+2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com>
+
+ * c-parser.cc (c_parser_omp_clause_from_to): Parse 'iterator' modifier.
+ * c-typeck.cc (c_finish_omp_clauses): Finish iterators for to/from
+ clauses.
+
2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com>
* c-parser.cc (c_parser_omp_clause_map): Parse 'iterator' modifier.
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index e0aafc405545..55636680718e 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -19944,8 +19944,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,
@@ -19957,20 +19960,27 @@ c_parser_omp_clause_from_to (c_parser *parser, enum
omp_clause_code kind,
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_COMMA)
- pos += 2;
- else if (c_parser_peek_nth_token_raw (parser, pos + 1)->type
- == CPP_OPEN_PAREN)
+ const char *identifier =
+ IDENTIFIER_POINTER (c_parser_peek_nth_token_raw (parser, pos)->value);
+ if (c_parser_peek_nth_token_raw (parser, pos + 1)->type
+ == CPP_OPEN_PAREN)
{
unsigned int npos = pos + 2;
if (c_parser_check_balanced_raw_token_sequence (parser, &npos)
- && (c_parser_peek_nth_token_raw (parser, npos)->type
- == CPP_CLOSE_PAREN))
- pos = npos + 1;
+ && (c_parser_peek_nth_token_raw (parser, npos)->type
+ == CPP_CLOSE_PAREN))
+ {
+ if (strcmp (identifier, "iterator") == 0)
+ iterator_length = npos - pos + 1;
+ pos = npos;
+ }
}
+ 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)
@@ -19983,6 +19993,7 @@ c_parser_omp_clause_from_to (c_parser *parser, enum
omp_clause_code kind,
int present_modifier = false;
int mapper_modifier = false;
tree mapper_name = NULL_TREE;
+ tree iterators = NULL_TREE;
for (int pos = 1; pos < colon_pos; ++pos)
{
@@ -20004,6 +20015,17 @@ c_parser_omp_clause_from_to (c_parser *parser, enum
omp_clause_code kind,
present_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);
@@ -20060,7 +20082,7 @@ c_parser_omp_clause_from_to (c_parser *parser, enum
omp_clause_code kind,
else
{
c_parser_error (parser, "%<to%> or %<from%> clause with modifier "
- "other than %<present%> or %<mapper%>");
+ "other than %<iterator%>, %<mapper%> or %<present%>");
parens.skip_until_found_close (parser);
return list;
}
@@ -20097,6 +20119,19 @@ c_parser_omp_clause_from_to (c_parser *parser, enum
omp_clause_code kind,
OMP_CLAUSE_CHAIN (last_new) = name;
}
+ 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_ITERATORS (c) = iterators;
+
return nl;
}
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index c75322e732e3..e6ec927c6bc2 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -16192,6 +16192,9 @@ c_finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE)
break;
+ /* FALLTHRU */
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
if (OMP_CLAUSE_ITERATORS (c)
&& c_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c)))
{
@@ -16199,8 +16202,6 @@ c_finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
break;
}
/* FALLTHRU */
- case OMP_CLAUSE_TO:
- case OMP_CLAUSE_FROM:
case OMP_CLAUSE__CACHE_:
{
using namespace omp_addr_tokenizer;
diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index cbc4eb230486..80d2c3b5069e 100644
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,9 @@
+2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com>
+
+ * parser.cc (cp_parser_omp_clause_from_to): Parse 'iterator' modifier.
+ * semantics.cc (finish_omp_clauses): Finish iterators for to/from
+ clauses.
+
2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com>
* parser.cc (cp_parser_omp_clause_map): Parse 'iterator' modifier.
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 1e9b7469b743..a13b0a9dc025 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -42417,8 +42417,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,
@@ -42429,14 +42432,25 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum
omp_clause_code kind,
int pos = 1;
int colon_pos = 0;
+ int iterator_length = 0;
while (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_NAME)
{
+ const char *identifier =
+ IDENTIFIER_POINTER (cp_lexer_peek_nth_token (parser->lexer,
+ pos)->u.value);
+ if (cp_lexer_nth_token_is (parser->lexer, pos + 1, CPP_OPEN_PAREN))
+ {
+ int n = cp_parser_skip_balanced_tokens (parser, pos + 1);
+ if (n != pos + 1)
+ {
+ if (strcmp (identifier, "iterator") == 0)
+ iterator_length = n - pos;
+ pos = n - 1;
+ }
+ }
if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)
pos += 2;
- else if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type
- == CPP_OPEN_PAREN)
- pos = cp_parser_skip_balanced_tokens (parser, pos + 1);
else
pos++;
if (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_COLON)
@@ -42449,6 +42463,7 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum
omp_clause_code kind,
bool present_modifier = false;
bool mapper_modifier = false;
tree mapper_name = NULL_TREE;
+ tree iterators = NULL_TREE;
for (int pos = 1; pos < colon_pos; ++pos)
{
@@ -42473,6 +42488,21 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum
omp_clause_code kind,
present_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);
@@ -42535,7 +42565,8 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum
omp_clause_code kind,
else
{
cp_parser_error (parser, "%<to%> or %<from%> clause with "
- "modifier other than %<present%> or %<mapper%>");
+ "modifier other than %<iterator%>, "
+ "%<mapper%> or %<present%>");
cp_parser_skip_to_closing_parenthesis (parser,
/*recovering=*/true,
/*or_comma=*/false,
@@ -42574,6 +42605,19 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum
omp_clause_code kind,
OMP_CLAUSE_CHAIN (last_new) = name;
}
+ 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_ITERATORS (c) = iterators;
+
return nl;
}
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 42a61574e9d2..dc102f15cf7c 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -8956,6 +8956,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type
ort)
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE)
break;
+ /* FALLTHRU */
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
if (OMP_CLAUSE_ITERATORS (c)
&& cp_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c)))
{
@@ -8963,8 +8966,6 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type
ort)
break;
}
/* FALLTHRU */
- case OMP_CLAUSE_TO:
- case OMP_CLAUSE_FROM:
case OMP_CLAUSE__CACHE_:
{
using namespace omp_addr_tokenizer;
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 78678c14b93e..cd687fc9f6ef 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -13644,7 +13644,8 @@ omp_instantiate_implicit_mappers (splay_tree_node n,
void *data)
static void
gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
enum omp_region_type region_type,
- enum tree_code code)
+ enum tree_code code,
+ gimple_seq *loops_seq_p = NULL)
{
using namespace omp_addr_tokenizer;
struct gimplify_omp_ctx *ctx, *outer_ctx;
@@ -14412,23 +14413,24 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
: TYPE_SIZE_UNIT (TREE_TYPE (decl));
- if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
- NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
+ 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)
{
remove = true;
+ exit_omp_iterator_loop_context (c);
break;
}
if (!DECL_P (decl))
{
- if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p,
- NULL, is_gimple_lvalue, fb_lvalue)
- == GS_ERROR)
- {
- remove = true;
- break;
- }
+ if (gimplify_expr (&OMP_CLAUSE_DECL (c), seq_p, NULL,
+ is_gimple_lvalue, fb_lvalue) == GS_ERROR)
+ remove = true;
+ exit_omp_iterator_loop_context (c);
break;
}
+ exit_omp_iterator_loop_context (c);
goto do_notice;
case OMP_CLAUSE__MAPPER_BINDING_:
@@ -19156,7 +19158,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
if ((ort & ORT_ACC) == 0)
in_omp_construct = false;
gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
- TREE_CODE (expr));
+ TREE_CODE (expr), &iterator_loops_seq);
if (TREE_CODE (expr) == OMP_TARGET)
optimize_target_teams (expr, pre_p);
if ((ort & (ORT_TARGET | ORT_TARGET_DATA)) != 0
@@ -19404,10 +19406,16 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq
*pre_p)
default:
gcc_unreachable ();
}
+
+ gimple_seq iterator_loops_seq = NULL;
+ remove_unused_omp_iterator_vars (&OMP_STANDALONE_CLAUSES (expr));
+ build_omp_iterators_loops (&OMP_STANDALONE_CLAUSES (expr),
+ &iterator_loops_seq);
+
gimplify_scan_omp_clauses (&OMP_STANDALONE_CLAUSES (expr), pre_p,
- ort, TREE_CODE (expr));
+ ort, TREE_CODE (expr), &iterator_loops_seq);
gimplify_adjust_omp_clauses (pre_p, NULL, &OMP_STANDALONE_CLAUSES (expr),
- TREE_CODE (expr));
+ TREE_CODE (expr), &iterator_loops_seq);
if (TREE_CODE (expr) == OACC_UPDATE
&& omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
OMP_CLAUSE_IF_PRESENT))
@@ -19471,7 +19479,8 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq
*pre_p)
gcc_unreachable ();
}
}
- stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
+ stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr),
+ iterator_loops_seq);
gimplify_seq_add_stmt (pre_p, stmt);
*expr_p = NULL_TREE;
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index bb5456a70eef..55b0ec12dcb3 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,9 @@
+2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com>
+
+ * c-c++-common/gomp/target-update-iterators-1.c: New.
+ * c-c++-common/gomp/target-update-iterators-2.c: New.
+ * c-c++-common/gomp/target-update-iterators-3.c: New.
+
2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com>
* c-c++-common/gomp/map-6.c (foo): Amend expected error message.
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-1.c
b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-1.c
new file mode 100644
index 000000000000..64602d45d494
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-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. or .from. clause with modifier other than .iterator., .mapper.
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-iterators-2.c
b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-2.c
new file mode 100644
index 000000000000..ae0a222485a0
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-2.c
@@ -0,0 +1,23 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+void f (int *x, float *y, double *z)
+{
+ #pragma omp target update to(iterator(i=0:10): x) /* { dg-warning "iterator
variable .i. not used in clause expression" }*/
+ ;
+
+ #pragma omp target update from(iterator(i2=0:10, j2=0:20): x[i2]) /* {
dg-warning "iterator variable .j2. not used in clause expression" }*/
+ ;
+
+ #pragma omp target update to(iterator(i3=0:10, j3=0:20, k3=0:30): x[i3+j3],
y[j3+k3], z[k3+i3])
+ /* { dg-warning "iterator variable .i3. not used in clause expression" "" {
target *-*-* } .-1 } */
+ /* { dg-warning "iterator variable .j3. not used in clause expression" "" {
target *-*-* } .-2 } */
+ /* { dg-warning "iterator variable .k3. not used in clause expression" "" {
target *-*-* } .-3 } */
+ ;
+}
+
+/* { dg-final { scan-tree-dump "update to\\\(x " "gimple" } } */
+/* { dg-final { scan-tree-dump "update from\\\(iterator\\\(int i2=0:10:1,
loop_label=" "gimple" } } */
+/* { dg-final { scan-tree-dump "to\\\(iterator\\\(int i3=0:10:1, int
k3=0:30:1, loop_label=" "gimple" } } */
+/* { dg-final { scan-tree-dump "to\\\(iterator\\\(int j3=0:20:1, int
k3=0:30:1, loop_label=" "gimple" } } */
+/* { dg-final { scan-tree-dump "to\\\(iterator\\\(int i3=0:10:1, int
j3=0:20:1, loop_label=" "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c
b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c
new file mode 100644
index 000000000000..ef55216876f6
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-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, loop_label=<D\.\[0-9\]+>, elems=omp_iter_data\.\[0-9\]+,
index=D\.\[0-9\]+\\):\\*D\.\[0-9\]+" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "from\\(iterator\\(int i=0:10:1,
loop_label=<D\.\[0-9\]+>, elems=omp_iter_data\.\[0-9\]+,
index=D\.\[0-9\]+\\):\\*D\.\[0-9\]+" 1 "gimple" } } */
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 5a0c4fc43fc7..cbdcd3bc0a1a 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1270,6 +1270,11 @@ 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:");
+ if (OMP_CLAUSE_ITERATORS (clause))
+ {
+ dump_omp_iterators (pp, OMP_CLAUSE_ITERATORS (clause), spc, flags);
+ pp_colon (pp);
+ }
dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
spc, flags, false);
goto print_clause_size;
@@ -1278,6 +1283,11 @@ dump_omp_clause (pretty_printer *pp, tree clause, int
spc, dump_flags_t flags)
pp_string (pp, "to(");
if (OMP_CLAUSE_MOTION_PRESENT (clause))
pp_string (pp, "present:");
+ if (OMP_CLAUSE_ITERATORS (clause))
+ {
+ dump_omp_iterators (pp, OMP_CLAUSE_ITERATORS (clause), spc, flags);
+ pp_colon (pp);
+ }
dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
spc, flags, false);
goto print_clause_size;
diff --git a/gcc/tree.cc b/gcc/tree.cc
index 5de1c3aa3815..d3dfd8b1d507 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -264,8 +264,8 @@ unsigned const char omp_clause_num_ops[] =
1, /* OMP_CLAUSE_IS_DEVICE_PTR */
1, /* OMP_CLAUSE_INCLUSIVE */
1, /* OMP_CLAUSE_EXCLUSIVE */
- 2, /* OMP_CLAUSE_FROM */
- 2, /* OMP_CLAUSE_TO */
+ 3, /* OMP_CLAUSE_FROM */
+ 3, /* OMP_CLAUSE_TO */
3, /* OMP_CLAUSE_MAP */
1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */
1, /* OMP_CLAUSE_DOACROSS */
diff --git a/gcc/tree.h b/gcc/tree.h
index 5dec4c0a8c62..e6268c1eb076 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1635,11 +1635,13 @@ class auto_suppress_location_wrappers
#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_CODE (NODE) == OMP_CLAUSE_FROM \
+ || OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_TO \
+ || 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_FROM, \
OMP_CLAUSE_MAP), 2)
/* True on OMP_FOR and other OpenMP/OpenACC looping constructs if the loop nest
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index c93ae3b4fd5e..ca16b1e11a0b 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,11 @@
+2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com>
+
+ * 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.
+
2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com>
* target.c (kind_to_name): New.
diff --git a/libgomp/target.c b/libgomp/target.c
index c0532ba1a673..aadb43731711 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2449,6 +2449,8 @@ 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;
+ size_t *iterator_count = NULL;
if (!devicep)
return;
@@ -2456,6 +2458,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, &iterator_count);
+
gomp_mutex_lock (&devicep->lock);
if (devicep->state == GOMP_DEVICE_FINALIZED)
{
@@ -2593,6 +2599,14 @@ 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);
+ free (iterator_count);
+ }
}
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 000000000000..5a4cad5c2195
--- /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 000000000000..93438d01c97e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c
@@ -0,0 +1,58 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* 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 000000000000..a70b21c4b75a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c
@@ -0,0 +1,67 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* 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;
+}
