Hi Tobias,
On 14/09/2026 12:13, Tobias Burnus wrote:
Hi PA,
thanks for the update.
Paul-Antoine Arras wrote:
I added usm_env_handling-4.c and usm_env_handling-4-self.c. Both pass
on my local GPU (nvptx). However, on gfx90a, all usm_env_handling-?.c
tests (the ones setting GOMP_RUNTIME_USM "enabled") show up as
unsupported. Is this expected?
[This was seemingly due DejaGNU not supporting dg-set-target-env-var with
remote testing; ]
I ran the tests manually, all pass.
Anyway, this was good to add this test as it uncovered missing
conditions for GOMP_MAP_TO_GRID and GOMP_MAP_FROM_GRID in
scan_sharing_clauses.
:-)
* * *
LGTM.
Thanks,
Tobias
PS: 'interdiff' looks scary, until one realizes that it failed half-way
through.
directly comparing works (graphical 'diff'), but the difference also
contains all
the post 1/3 changes that were reviewed before.
(Hopefully, this will get better when we switch to the forge.)
* * *
While going through the Fortran review, I noticed that there was a case
working in C and not in C++:
#pragma omp target update to(arr[::2])
The literal 2 would get rejected because it was treated as an
id-expression. This is now fixed in 0001-WIP-C-fixup-for-non-id-expr.patch.
* * *
Something that also came up as part of the Fortran review, but actually
belongs in the middle end, is the lack of a positive check for a stride
that is not known at compile time. This check is now added in omp-low.cc
and fires at run time whenever the stride is not a compile-time
constant. See
0002-ME-fixup-to-check-at-runtime-for-non-positive-stride.patch.
Thanks,
--
PA
From 4f093fe9b4aa3083995b66de263cc29524961315 Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <[email protected]>
Date: Wed, 16 Sep 2026 17:23:17 +0200
Subject: [PATCH 1/2] WIP C++ fixup for '[::<non-id-expr>'
---
gcc/cp/parser.cc | 33 ++++++++++++++-----
.../c-c++-common/gomp/array-section-1.c | 4 +--
gcc/testsuite/g++.dg/gomp/array-section-7.C | 18 ++++++++++
3 files changed, 43 insertions(+), 12 deletions(-)
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 58e6810b710..b16ff026d55 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -9378,6 +9378,22 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
bool double_colon_only
= cp_lexer_next_token_is (parser->lexer, CPP_SCOPE)
&& cp_lexer_nth_token_is (parser->lexer, 2, CPP_CLOSE_SQUARE);
+ bool triple_colon = cp_lexer_next_token_is (parser->lexer, CPP_SCOPE)
+ && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON);
+
+ bool double_colon_stride = false;
+ if (parser->omp_array_section_kind == OMP_ARRAY_SECTION_STRIDED
+ && cp_lexer_next_token_is (parser->lexer, CPP_SCOPE) && !double_scope
+ && !double_colon_only && !triple_colon)
+ {
+ cp_parser_parse_tentatively (parser);
+ cp_parser_expression (parser, NULL, /*cast_p=*/false,
+ /*decltype_p=*/false,
+ /*warn_comma_p=*/warn_comma_subscript);
+ if (cp_parser_error_occurred (parser))
+ double_colon_stride = true;
+ cp_parser_abort_tentative_parse (parser);
+ }
saved_greater_than_is_operator_p = parser->greater_than_is_operator_p;
parser->greater_than_is_operator_p = true;
@@ -9400,9 +9416,9 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
if (for_offsetof)
index = cp_parser_constant_expression (parser);
else if (parser->omp_array_section_kind == OMP_ARRAY_SECTION_STRIDED
- && (double_scope || double_colon_only))
+ && (double_scope || double_colon_only || double_colon_stride))
{
- /* Handle '[::::' and '[::]' in post_colon_parsing. */
+ /* Handle '[::::', '[::]' and '[::<non-id-expr>' in post_colon_parsing. */
}
else if (parser->omp_array_section_kind == OMP_ARRAY_SECTION_NONE
|| (cp_lexer_next_token_is_not (parser->lexer, CPP_COLON)
@@ -9467,8 +9483,7 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
maybe_warn_cpp0x (CPP0X_INITIALIZER_LISTS);
index = cp_parser_braced_list (parser);
}
- else if (cp_lexer_next_token_is (parser->lexer, CPP_SCOPE)
- && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+ else if (triple_colon)
{
/* The lexer greedily tokenizes '[:::' as '[ :: :'.
It should be '[ : ::'. Let's swap the last two. */
@@ -9504,12 +9519,12 @@ post_colon_parsing:
&& (open_splice || cp_lexer_next_token_is (parser->lexer, CPP_COLON)
|| cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_SPLICE)
|| (parser->omp_array_section_kind == OMP_ARRAY_SECTION_STRIDED
- && (double_scope || double_colon_only))))
+ && (double_scope || double_colon_only || double_colon_stride))))
{
tree length = NULL_TREE, stride = NULL_TREE;
bool close_splice = cp_lexer_next_token_is (parser->lexer,
CPP_CLOSE_SPLICE);
- if (!open_splice && !double_scope)
+ if (!open_splice && !double_scope && !double_colon_stride)
cp_lexer_consume_token (parser->lexer);
if (open_splice && close_splice)
{
@@ -9527,7 +9542,7 @@ post_colon_parsing:
cp_lexer_peek_nth_token (parser->lexer, 2)->type = CPP_SCOPE;
}
else if (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_SQUARE)
- && !close_splice && !double_scope
+ && !close_splice && !double_scope && !double_colon_stride
&& cp_lexer_next_token_is_not (parser->lexer, CPP_COLON))
{
if (cxx_dialect >= cxx23)
@@ -9561,8 +9576,8 @@ post_colon_parsing:
}
if (parser->omp_array_section_kind == OMP_ARRAY_SECTION_STRIDED
- && (cp_lexer_next_token_is (parser->lexer, CPP_COLON)
- || double_scope))
+ && (cp_lexer_next_token_is (parser->lexer, CPP_COLON) || double_scope
+ || double_colon_stride))
{
cp_lexer_consume_token (parser->lexer);
/* We could check for C++-23 multidimensional/comma-separated
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-1.c b/gcc/testsuite/c-c++-common/gomp/array-section-1.c
index c0561284ac4..5af3f950dd7 100644
--- a/gcc/testsuite/c-c++-common/gomp/array-section-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-1.c
@@ -1,14 +1,12 @@
/* { dg-do compile } */
-/* Check that array section syntax with empty length is handled correctly.
- '::1' is valid only in C. */
+/* Check that array section syntax with empty length is handled correctly. */
int n;
void f() {
int A[10], B[10][10];
#pragma omp target update to(A[::], B[::][::1])
-/* { dg-error "expected id-expression before numeric constant" "" { target c++ } .-1 } */
#pragma omp target update to(A[::], B[0][: : ])
#pragma omp target update to(A[1:2:], B[:1:][1: : ])
diff --git a/gcc/testsuite/g++.dg/gomp/array-section-7.C b/gcc/testsuite/g++.dg/gomp/array-section-7.C
index 5dc87030c08..9fb282bfde4 100644
--- a/gcc/testsuite/g++.dg/gomp/array-section-7.C
+++ b/gcc/testsuite/g++.dg/gomp/array-section-7.C
@@ -3,6 +3,9 @@
// Check parsing of double colon (::). In case of ambiguity between the scope
// resolution operator and the array section syntax, spaces are required.
+// There is only an actual ambiguity when what follows "::" could itself be
+// a qualified-id (e.g. an identifier); for anything else (a literal, a
+// parenthesized expression, ...) no space is needed.
int lb2 = 4;
int len2 = 16;
@@ -29,6 +32,21 @@ void g() {
#pragma omp target update to(arr[::])
// { dg-final { scan-tree-dump "to\\(arr\\\[0\\\] \\\[len: 512\\\]\\)" "original" } }
+ // Valid: no ambiguity, since a numeric literal can never start an
+ // id-expression, so "2" is unambiguously the stride.
+#pragma omp target update to(arr[::2])
+// { dg-final { scan-tree-dump "map\\(to_grid:arr \\\[len: 4\\\]\\) map\\(grid_dim:0 \\\[len: 64\\\]\\) map\\(grid_stride:2\\)" "original" } }
+
+ // Likewise for a parenthesized expression as the stride.
+#pragma omp target update to(arr[::(s)])
+// { dg-final { scan-tree-dump "map\\(to_grid:arr \\\[len: 4\\\]\\) map\\(grid_dim:0 \\\[len: 128 /\\\[cl\\\] \\(sizetype\\) s\\\]\\) map\\(grid_stride:s\\)" "original" } }
+
+ // A negative or zero stride is rejected as a semantic error (not a parse
+ // failure), same as when written with a disambiguating space.
+#pragma omp target update to(arr[::-2])
+// { dg-error "negative or zero stride" "" { target *-*-* } .-1 }
+// { dg-error "must contain at least one" "" { target *-*-* } .-2 }
+
// Valid: len2 is looked up in the global namespace and forms the length of the
// array section.
#pragma omp target update to(arr[:::len2])
--
2.55.0
From 0cbafd6f8adf044fd93b665552fd15849efa5594 Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <[email protected]>
Date: Wed, 16 Sep 2026 20:15:22 +0200
Subject: [PATCH 2/2] ME fixup to check at runtime for non-positive stride
---
gcc/omp-low.cc | 45 +++++++++++++++++++
.../target-update-negative-stride-1.c | 26 +++++++++++
2 files changed, 71 insertions(+)
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-update-negative-stride-1.c
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 383eb8dd9c1..819c3a0b149 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -13738,6 +13738,51 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
== GOMP_MAP_GRID_STRIDE))
{
stride = OMP_CLAUSE_DECL (nc2);
+
+ if (TREE_CODE (stride) != INTEGER_CST)
+ {
+ /* Check at run time that stride is positive. */
+ location_t loc = gimple_location_safe (ilist);
+ tree err
+ = builtin_decl_explicit (BUILT_IN_GOMP_ERROR);
+ const char *str
+ = "Non-positive stride in array section";
+ tree msg
+ = build_string_literal (strlen (str) + 1,
+ str);
+ tree msglen
+ = build_int_cst (size_type_node, -2);
+ err = build_call_expr_loc (loc, err, 2, msg,
+ msglen);
+ tree l1
+ = create_artificial_label (UNKNOWN_LOCATION);
+ tree l2
+ = create_artificial_label (UNKNOWN_LOCATION);
+ tree l3
+ = create_artificial_label (UNKNOWN_LOCATION);
+
+ gimple_seq new_body = NULL;
+
+ gcond *cond = gimple_build_cond (
+ LE_EXPR, stride,
+ build_int_cst (TREE_TYPE (stride), 0), l1,
+ l2);
+ gimple_seq_add_stmt (&new_body, cond);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_label (l1));
+ gimplify_and_add (err, &new_body);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_goto (l3));
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_label (l2));
+ gimple_seq_add_seq (&new_body, ilist);
+ gimple_seq_add_stmt (&new_body,
+ gimple_build_label (l3));
+ ilist = gimple_build_bind (NULL_TREE, new_body,
+ make_node (BLOCK));
+ gimple_set_location (ilist, loc);
+ }
+
stride = fold_convert (sizetype, stride);
if (OMP_CLAUSE_SIZE (nc2))
{
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-negative-stride-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-negative-stride-1.c
new file mode 100644
index 00000000000..aeb0ff1ce54
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-negative-stride-1.c
@@ -0,0 +1,26 @@
+/* { dg-shouldfail "non-positive stride in array section" } */
+
+/* A stride that is only known to be non-positive at run time (as opposed to
+ a literal, which is rejected at compile time already) must still be
+ rejected, not silently ignored. */
+
+int
+main (void)
+{
+ int n = 20;
+ int arr[20];
+
+ for (int i = 0; i < n; i++)
+ arr[i] = i;
+
+ int stride = -2;
+
+ #pragma omp target data map(tofrom: arr[0:20])
+ {
+ #pragma omp target update from(arr[0:10:stride])
+ }
+
+ return 0;
+}
+
+/* { dg-output "libgomp: fatal error: Non-positive stride in array section" } */
--
2.55.0
