Hi Tobias,
Thanks for the review. I'm replying to both your emails about C++ here.
On 21/08/2026 11:31, Tobias Burnus wrote:
Hi PA,
some first remarks; I still need to read the FE bits and come
back to some parts I was wondering about - and that I wanted to
test via the attached testcase.
In particular, I believe there are some issues:
* I think stride should be stride (offset) multiplier
(see below) - at least for Fortran, I think that's
needed - if the code I have below is valid.
* I have a parse error - and, IMGO, a bogus error message
* The IMHO valid testcase crashes - at least with nvptx
offloading - in omp_target_memcpy_rect_worker
(or actually in cuMemcpy3D_v2 invoked via GOMP_OFFLOAD_memcpy3d)
Can you check?
* * *
BTW: Possibly after the last patch in the series, we also want to
update https://gcc.gnu.org/onlinedocs/libgomp/OpenMP-5_002e0.html
From u'N'implemented to 'P'artial or even, 'Y'es - fully for the
three items: array shaping and array-sections with strides ('Y'
once C++ and C patches have landed) - and for
'Discontiguous array section with target update' either 'P'
(plus note) or 'Y'.
* * *
Paul-Antoine Arras wrote:
From: Julian Brown<[email protected]>
This patch adds support for OpenMP 5.0 strided updates and the
array-shaping operator ("([x][y][z]) foo[0:n]..."). This is mostly for
C++ only so far, though necessary changes have been made to the C FE to
adjust for changes to shared data structures.
I am wondering whether we need to take care of the following
issue that exists in Fortran:
type t2; ...; end type t2 !
type t
type(t2) :: x ! sizeof = 40
integer :: y ! sizeof = 4
end type t ! → sizeof = 44
type(t) :: arr(100)
type(t2) :: p(:)
p => arr(:)%x
Here, sizeof(p(1)) == 40
and loc(p(2)) - loc(p(1)) == 44.
And p2 => p(::2)
Here, the stride is syntax wise 2 but internally loc(p2(2)) - loc(p2(1))
= 2*44.
But that stride is not a multiple of sizeof(type(t2)) but of
sizeof(type(t)).
The solution on the Fortran side is to use a stride multiplier ('sm')
that is
a byte offset instead of a stride (with sm = stride * elem_size).
* * *
The question is whether we need to handle non-integer multiples
of the element size for strided updates? (And even if not, does
it make sense to be future proof?)
In particular:
* * *
Is the following valid?
!$omp target enter data map(arr) ! argument must be contiguous (*)
...
!$omp target update to(p(::)) ! or 'to(p(::2))'
The first one needs to be contiguous (quote from TR15):
"Additional restrictions to the map clause for mapping of arrays,
array sections, and structures are as follows:
* If a list item is an array or an array section, it must specify
contiguous storage."
And the wording in "12 Data-Motion Control" implies that this is
valid:
"If a list item is an array or array section then it is treated as
if it is replaced by each of its array elements in the clause."
albeit the following is slightly confusing:
* * *
If so, the stride to stride multiplier conversion of libgomp/target.c's
omp_target_memcpy_rect_worker, i.e.
+ && __builtin_mul_overflow (element_size, strides[0], &stride))
should be moved to the compiler itself.
I think comments / variable names should make clear whether they
refers to a multiple of the array-element size or a (stride)
offset/stride multiplier.
* * *
Syntax wise, Fortran permits 'var(:)%x' or 'var(::2)%x'; I think
OpenMP only permits an array section (w/ and w/o strides) when it
is the last designator, i.e. the issue only occurs when doing the
detour via either a pointer or via:
associate(p => arr(:)%x)
... to(p(:)) ! note: 'tp(p)' is invalid as not contiguous ... (**)
end associate
where associate is an alias, if the RHS counts as variable, then
it is also a variable, otherwise only the result of an expression.
(**) Namely, "12 Data-Motion Control" requires:
"* If an array appears as a list item in a data-motion clause and
it has corresponding storage in the device data environment, the
corresponding storage must correspond to a single mappable storage
block that was previously mapped.
• If a list item in a data-motion clause has corresponding storage
in the device data environment, all corresponding storage must
correspond to a single mappable storage block that was previously
mapped."
where the first one seems to apply to 'p' and the second one to
'p(:)' or 'p(4:6)' or 'p(::2)'.
I believe this Fortran-specific issue is addressed in the Fortran patch
at the end of this series (3/3). If turns out not to be, I'm happy to
rediscuss it as part of that patch's review.
* * *
TESTCASE
========
I have tried the attached testcase - but it fails to compile with
two errors, which look bogus to me:
int x[10][8][6];
...
#pragma omp target update from(x[: :4][: :3][: :2], ...)
foo.C:32:42: error: expected primary-expression before »:« token
32 | #pragma omp target update from(x[: :4][: :3][: :2], ...
| ^
And when inserting '10':
foo.C:33:66: error: length '10' with stride '4' above array section size
in 'from' clause
33 | #pragma omp target update from(x[: 10 :4][: 8 :3][:
6 :2], ...
| ^
But I expect that I can access element 0, 4 and 8 using the [0:10:4]
syntax?
That's correct.
Your testcase actually highlighted several issues in this patch:
* bogus parse error;
* missing default for length;
* missing dereference in array-shape cast.
These are fixed in the attached fixup patch and exercised by new test
libgomp/testsuite/libgomp.c++/array-shaping-14.C.
Replacing the literals by variables (int four = 4 etc.) makes it compile
but
then it fails at runtime as mentioned above.
* * *
2026-08-13 Paul-Antoine Arras
gcc/c-family/
* c-omp.cc (omp_handle_noncontig_array): Propagate
OMP_CLAUSE_ITERATORS onto synthesised grid clauses.
gcc/cp/
* parser.cc (cp_parser_postfix_open_square_expression): Only
accept a 2nd colon (stride) for to/from clauses, not map.
(cp_parser_omp_var_list_no_open): Only count OMP_ARRAY_SECTION
layers, not every ARRAY_REF layer, in a to/from clause's base
decl.
* parser.h (cp_parser::omp_array_section_p): Replace bool with...
(cp_parser::omp_array_section_kind): ...this 3-state unsigned char.
(OMP_ARRAY_SECTION_NONE, OMP_ARRAY_SECTION_UNSTRIDED,
OMP_ARRAY_SECTION_STRIDED): Define.
* semantics.cc (handle_omp_array_sections): Handle a genuinely
discontiguous update at any access-chain depth.
(finish_omp_clauses): Splice the GOMP_MAP_TO_GRID/FROM_GRID
replacement clause in at *PC when DISCONTIGUOUS == 2.
gcc/
* gimplify.cc (gimplify_scan_omp_clauses): Handle
OMP_TARGET_UPDATE like the other target constructs.
(gimplify_adjust_omp_clauses): Fix a leaked gimplify context in
the GOMP_MAP_GRID_DIM/GRID_STRIDE case.
gimplify.cc (remove_unused_omp_iterator_vars): Exempt
GOMP_MAP_GRID_DIM/GRID_STRIDE clauses.
(gimplify_adjust_omp_clauses): Restrict the "unsupported map
expression" sorry to GOMP_MAP_GRID_DIM/GRID_STRIDE clauses.
(gimplify_omp_target_update): Reject an iterator combined with a
grid map clause.
I was wondering whether a testcase needs to be added for those?
(Including for C++26?)
The (disallowed, for now) interaction between map iterators and strided
update is tested by c-c++-common/gomp/target-update-iterators-4.c.
Discontiguous update at any access-chain depth is covered by
array-shaping-8.C/-9.C.
Multi-item update directives are checked in array-shaping-14.C, added to
the fixup patch.
I also just added bad-array-section-12.C to check that a stride is
rejected for a map clause but accepted for update to/from.
As regards C++26, I folded into the fixup patch, the patch "C++:
Disambiguate colon syntaxes in array sections" submitted to OG16:
https://gcc.gnu.org/pipermail/gcc-patches/2026-August/727096.html.
* * *
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -11046,6 +11049,19 @@ omp_group_last (tree *start_p)
grp_last_p = &OMP_CLAUSE_CHAIN (c);
break;
+ case GOMP_MAP_TO_GRID:
+ case GOMP_MAP_FROM_GRID:
+ while (nc
+ && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE))
+ {
+ grp_last_p = &OMP_CLAUSE_CHAIN (c);
+ c = nc;
+ nc = OMP_CLAUSE_CHAIN (c);
+ }
Indentation issue.
I don't see such issue on the code itself, maybe just a diff quirk?
case GOMP_MAP_TO_GRID:
case GOMP_MAP_FROM_GRID:
while (nc
&& OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM
|| OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE))
{
grp_last_p = &OMP_CLAUSE_CHAIN (c);
c = nc;
nc = OMP_CLAUSE_CHAIN (c);
}
break;
* * *
+ /* The OMP_CLAUSE_DECL for GRID_DIM/GRID_STRIDE isn't
necessarily
+ an lvalue -- e.g. it might be a constant. So handle it
+ specially here. */
[Off topic: OpenMP meanwhile also permits constexpr, const etc. in map
clauses – which the implementation may/has to ignore. But I guess we
want to handle them already in the front end - such that this comment
can stay.]
* * *
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1423,7 +1423,7 @@ DEFTREECODE (OMP_ATOMIC_CAPTURE_NEW,
"omp_atomic_capture_new", tcc_statement, 2)
DEFTREECODE (OMP_CLAUSE, "omp_clause", tcc_exceptional, 0)
/* An OpenMP array section. */
-DEFTREECODE (OMP_ARRAY_SECTION, "omp_array_section", tcc_expression, 3)
+DEFTREECODE (OMP_ARRAY_SECTION, "omp_array_section", tcc_expression, 4)
I wonder whether we should expand this, e.g. append "with strides" or
maybe even better "(decl, low bound, length, stride)" - such that one
knows what's actually stored in the four operands.
Expanded the comment with a description for each operand.
* * *
I have tried the attached testcase - but it fails to compile with
two errors, which look bogus to me:
int x[10][8][6];
...
#pragma omp target update from(x[: :4][: :3][: :2], ...)
foo.C:32:42: error: expected primary-expression before »:« token
32 | #pragma omp target update from(x[: :4][: :3][: :2], ...
| ^
OpenMP specifies that if length is absent, it should be:
⌈(size − lower-bound)/stride⌉
(note the round up)
The default length is now computed and set in omp_expand_grid_dim.
And when inserting '10':
foo.C:33:66: error: length '10' with stride '4' above array section size in
'from' clause
33 | #pragma omp target update from(x[: 10 :4][: 8 :3][: 6 :2], ...
| ^
OK – C/C++ defines this differently than Fortran - in Fortran, it is
a bit like 'for (i = lower; i <= upper; i += stride)' but in OpenMP
it is for C/C++:
{ lower-bound,
lower-bound + stride,
lower-bound + 2 * stride,
... ,
lower-bound + ((length - 1) * stride) }
Thus, replace in my example:
// Still valid and still a parse error:
#pragma omp target update from(x[: :4][: :3][: :2], (([10][8][6])p)[: :4][:
:3][: :2])
// WRONG:
// #pragma omp target update from(x[: 10 :4][: 8 :3][: 6 :2],
(([10][8][6])p)[: 10 :4][: 8 :3][: 6 :2])
// SHOULD BE
#pragma omp target update from(x[: (10+3)/4 :4][: (8+2)/3 :3][: (6+1)/2 :2],
(([10][8][6])p)[: (10+3)/4 :4][: (8+2)/3 :3][: (6+1)/2 :2])
Testcase adapted into array-shaping-14.C.
* * *
However, I still run into the same run-time problem:
SEGFAULT in cuMemcpy3D_v2
... but the values seem to be fine at a glance:
#0 omp_target_memcpy_rect_worker (dst=0x7fffffffc4e8, src=0x7fffd97009f8, element_size=4,
num_dims=num_dims@entry=3, volume=0x4ad1d0 <main::.omp_len.31>, strides=0x4ad1f0
<main::.omp_stride.32>,
dst_offsets=0x5e0280 <main::.omp_index.30>, src_offsets=0x5e0280 <main::.omp_index.30>,
dst_dimensions=0x4ad1b0 <main::.omp_dim.29>, src_dimensions=0x4ad1b0 <main::.omp_dim.29>,
dst_devicep=0x0, src_devicep=0x627010,
tmp_size=0x7fffffffc440, tmp=0x7fffffffc448) at
/home/tob/repos/gcc/libgomp/target.c:5283
5283 {
(gdb) p volume[0]
$17 = 3
(gdb) p volume[1]
$18 = 3
(gdb) p volume[2]
$19 = 3
(gdb) p stride[0]
No symbol "stride" in current context.
(gdb) p strides[0]
$20 = 4
(gdb) p strides[1]
$21 = 3
(gdb) p strides[2]
$22 = 2
(gdb) p src_offset
No symbol "src_offset" in current context.
(gdb) p src_offsets
$23 = (const size_t *) 0x5e0280 <main::.omp_index.30>
(gdb) p src_offsets[0]
$24 = 0
(gdb) p src_offsets[1]
$25 = 0
(gdb) p src_offsets[2]
$26 = 0
(gdb) p dst_offsets[2]
$27 = 0
(gdb) p dst_offsets[1]
$28 = 0
(gdb) p dst_offsets[0]
$29 = 0
(gdb) p src_dimensions[0]
$30 = 10
(gdb) p src_dimensions[1]
$31 = 8
(gdb) p src_dimensions[2]
$32 = 6
(gdb) p dst_dimensions[2]
$33 = 6
(gdb) p dst_dimensions[1]
$34 = 8
(gdb) p dst_dimensions[0]
$35 = 10
This then calls:
GOMP_OFFLOAD_memcpy3d (dst_ord=-1, src_ord=0, dim2_size=12, dim1_len=3,
dim0_len=3, dst=0x7fffffffc4e8, dst_offset2_size=0, dst_offset1_len=0,
dst_offset0_len=0, dst_dim2_size=24, dst_dim1_len=8, src=0x7fffd97009f8,
src_offset2_size=0, src_offset1_len=0, src_offset0_len=0, src_dim2_size=24,
src_dim1_len=8) at /home/tob/repos/gcc/libgomp/plugin/plugin-nvptx.c:2512
2512 {
where I don't immediately see whether that's okay or not.
The error was a missing dereference in array-shape cast.
* * *
BTW: The following program fails with both the C and the C++ compiler:
void f()
{
int x[10];
int y[10][10][10];
#pragma omp target update to(x[0:0:-1]) // invalid - missing diagnostic:
stride must evaluate to a positive integer.
#pragma omp target update to(x[0:0:0]) // invalid - missing diagnostic:
stride must evaluate to a positive integer.
#pragma omp target update to(y[0:0:0]) // ICE (segfault)
}
Added adequate diagnostic and test bad-array-shaping-8.C. Also updated
bad-array-section-13.C.
Thanks,
--
PA
From 7278f18022621160f4725241fcdae880f57e9303 Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <[email protected]>
Date: Tue, 25 Aug 2026 17:22:58 +0200
Subject: [PATCH] OpenMP: Fix strided and shaped-array update issues for C++
This is a fixup commit for "OpenMP: Support strided and shaped-array
updates for C++", addressing issues found during review of that
patch.
gcc/c-family/ChangeLog:
* c-omp.cc (omp_expand_grid_dim): Add a TYPE output parameter.
Default an omitted low bound to 0, an omitted stride to 1, and
an omitted length to ceil((size - low bound) / stride).
(omp_handle_noncontig_array): Update call to omp_expand_grid_dim.
gcc/cp/ChangeLog:
* parser.cc (cp_parser_postfix_open_square_expression): Support
the C++26 scope (::) and splice ([:/:]) notations in
combination with strides in OpenMP array sections.
(cp_parser_omp_var_list_no_open): Treat an INDIRECT_REF base as
a bare pointer dereference when it is not already part of a reshape.
* semantics.cc (handle_omp_array_sections_1): Diagnose a
negative or zero stride.
(cp_build_omp_arrayshape_cast): Dereference a pointer argument
before reinterpreting it as the target array type.
gcc/ChangeLog:
* tree.def (OMP_ARRAY_SECTION): Document the four operands
(decl, low bound, length, stride).
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/target-update-iterators-4.c: Update expected
diagnostic wording.
* g++.dg/gomp/bad-array-shaping-8.C: Update expected diagnostic
for a negative-stride array section to the new, more specific
message.
* g++.dg/gomp/array-section-7.C: New test.
* g++.dg/gomp/array-section-8.C: New test.
* g++.dg/gomp/bad-array-section-12.C: New test.
* g++.dg/gomp/bad-array-section-13.C: New test.
libgomp/ChangeLog:
* testsuite/libgomp.c++/array-shaping-14.C: New test.
---
gcc/c-family/c-omp.cc | 40 +++++++--
gcc/cp/parser.cc | 88 +++++++++++++------
gcc/cp/semantics.cc | 24 +++++
.../gomp/target-update-iterators-4.c | 2 +-
gcc/testsuite/g++.dg/gomp/array-section-7.C | 58 ++++++++++++
gcc/testsuite/g++.dg/gomp/array-section-8.C | 36 ++++++++
.../g++.dg/gomp/bad-array-section-12.C | 32 +++++++
.../g++.dg/gomp/bad-array-section-13.C | 16 ++++
.../g++.dg/gomp/bad-array-shaping-8.C | 4 +-
gcc/tree.def | 6 +-
.../testsuite/libgomp.c++/array-shaping-14.C | 65 ++++++++++++++
11 files changed, 334 insertions(+), 37 deletions(-)
create mode 100644 gcc/testsuite/g++.dg/gomp/array-section-7.C
create mode 100644 gcc/testsuite/g++.dg/gomp/array-section-8.C
create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-section-12.C
create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-section-13.C
create mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-14.C
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 667936bc7d2..d46ac8476ef 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -3718,12 +3718,15 @@ omp_expand_access_chain (tree c, tree expr, vec<omp_addr_token *> &addr_tokens,
}
static tree *
-omp_expand_grid_dim (location_t loc, tree *pc, tree decl)
+omp_expand_grid_dim (location_t loc, tree *pc, tree decl, tree *type)
{
if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
- pc = omp_expand_grid_dim (loc, pc, TREE_OPERAND (decl, 0));
+ pc = omp_expand_grid_dim (loc, pc, TREE_OPERAND (decl, 0), type);
else
- return pc;
+ {
+ *type = TREE_TYPE (decl);
+ return pc;
+ }
tree c = *pc;
tree low_bound = TREE_OPERAND (decl, 1);
@@ -3732,10 +3735,35 @@ omp_expand_grid_dim (location_t loc, tree *pc, tree decl)
tree cd = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (cd, GOMP_MAP_GRID_DIM);
+
+ if (low_bound == NULL_TREE)
+ low_bound = size_zero_node;
OMP_CLAUSE_DECL (cd) = unshare_expr (low_bound);
+
+ if (stride == NULL_TREE)
+ stride = size_one_node;
+
+ if (length == NULL_TREE && TYPE_DOMAIN (*type))
+ {
+ /* size = maxval - minval + 1 */
+ tree dtype = TYPE_DOMAIN (*type);
+ tree minval = TYPE_MIN_VALUE (dtype);
+ tree maxval = TYPE_MAX_VALUE (dtype);
+ minval = fold_convert (sizetype, minval);
+ maxval = fold_convert (sizetype, maxval);
+ tree size = size_binop (MINUS_EXPR, maxval, minval);
+ size = size_binop (PLUS_EXPR, size, size_one_node);
+
+ /* length = ⌈(size − lower-bound)/stride⌉ */
+ low_bound = fold_convert (sizetype, low_bound);
+ length = size_binop (MINUS_EXPR, size, low_bound);
+ tree sstride = fold_convert (sizetype, stride);
+ length = size_binop (CEIL_DIV_EXPR, length, sstride);
+ }
+
OMP_CLAUSE_SIZE (cd) = unshare_expr (length);
- if (stride && !integer_onep (stride))
+ if (!integer_onep (stride))
{
tree cs = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (cs, GOMP_MAP_GRID_STRIDE);
@@ -3753,6 +3781,7 @@ omp_expand_grid_dim (location_t loc, tree *pc, tree decl)
pc = &OMP_CLAUSE_CHAIN (c);
}
+ *type = TREE_TYPE (*type);
return pc;
}
@@ -3790,7 +3819,8 @@ omp_handle_noncontig_array (location_t loc, tree *pc, tree c, tree base)
*pc = c_map;
- tree *ret = omp_expand_grid_dim (loc, pc, OMP_CLAUSE_DECL (c));
+ tree dtype;
+ tree *ret = omp_expand_grid_dim (loc, pc, OMP_CLAUSE_DECL (c), &dtype);
/* The GOMP_MAP_TO_GRID/FROM_GRID node above and the GOMP_MAP_GRID_DIM/
GRID_STRIDE nodes omp_expand_grid_dim just built all need the same
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 0f5c5d9e0b5..982422d9de8 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -9339,9 +9339,10 @@ cp_parser_parenthesized_expression_list_elt (cp_parser *parser, bool cast_p,
changes how we deal with integer constant expressions.
With parser->omp_array_section_kind, it also handles OpenMP
- array sections of the type [ index : length ] where both
- index and length are optional. Note that an absent index
- might be lexed as CPP_OPEN_SPLICE ('[:') since C++26. */
+ array sections of the type [ index : length : stride ] where index,
+ length and stride are optional. Note that an absent index
+ might be lexed as CPP_OPEN_SPLICE ('[:') and an absent stride
+ might be lexed as CPP_CLOSE_SPLICE (':]') since C++26. */
static tree
cp_parser_postfix_open_square_expression (cp_parser *parser,
@@ -9360,6 +9361,12 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
/* Consume the `[' token - or with open_splice the '[:' token. */
cp_lexer_consume_token (parser->lexer);
+ bool double_scope = cp_lexer_next_token_is (parser->lexer, CPP_SCOPE)
+ && cp_lexer_nth_token_is (parser->lexer, 2, CPP_SCOPE);
+ bool double_colon_only
+ = cp_lexer_next_token_is (parser->lexer, CPP_SCOPE)
+ && cp_lexer_nth_token_is (parser->lexer, 2, CPP_CLOSE_SQUARE);
+
saved_greater_than_is_operator_p = parser->greater_than_is_operator_p;
parser->greater_than_is_operator_p = true;
@@ -9380,6 +9387,11 @@ cp_parser_postfix_open_square_expression (cp_parser *parser,
constant expressions here. */
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))
+ {
+ /* Handle '[::::' and '[::]' in post_colon_parsing. */
+ }
else if (!parser->omp_array_section_kind
|| (cp_lexer_next_token_is_not (parser->lexer, CPP_COLON)
&& cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_SPLICE)))
@@ -9443,6 +9455,14 @@ 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))
+ {
+ /* The lexer greedily tokenizes '[:::' as '[ :: :'.
+ It should be '[ : ::'. Let's swap the last two. */
+ cp_lexer_peek_token (parser->lexer)->type = CPP_COLON;
+ cp_lexer_peek_nth_token (parser->lexer, 2)->type = CPP_SCOPE;
+ }
else
index = cp_parser_expression (parser, NULL, /*cast_p=*/false,
/*decltype_p=*/false,
@@ -9463,12 +9483,14 @@ post_colon_parsing:
if (parser->omp_array_section_kind
&& (open_splice || cp_lexer_next_token_is (parser->lexer, CPP_COLON)
- || cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_SPLICE)))
+ || cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_SPLICE)
+ || (parser->omp_array_section_kind == OMP_ARRAY_SECTION_STRIDED
+ && (double_scope || double_colon_only))))
{
tree length = NULL_TREE, stride = NULL_TREE;
bool close_splice = cp_lexer_next_token_is (parser->lexer,
CPP_CLOSE_SPLICE);
- if (!open_splice)
+ if (!open_splice && !double_scope)
cp_lexer_consume_token (parser->lexer);
if (open_splice && close_splice)
{
@@ -9476,8 +9498,17 @@ post_colon_parsing:
UNKNOWN_LOCATION);
length = error_mark_node;
}
+ else if (open_splice
+ && cp_lexer_next_token_is (parser->lexer, CPP_SCOPE)
+ && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+ {
+ /* The lexer greedily tokenizes '[::::' as '[: :: :'.
+ It should be '[ : : ::'. Let's swap the last two. */
+ cp_lexer_peek_token (parser->lexer)->type = CPP_COLON;
+ 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)
+ && !close_splice && !double_scope)
{
if (cxx_dialect >= cxx23)
{
@@ -9502,7 +9533,7 @@ post_colon_parsing:
length = error_mark_node;
}
}
- else
+ else if (cp_lexer_next_token_is_not (parser->lexer, CPP_COLON))
length
= cp_parser_expression (parser, NULL, /*cast_p=*/false,
/*decltype_p=*/false,
@@ -9510,7 +9541,8 @@ post_colon_parsing:
}
if (parser->omp_array_section_kind == OMP_ARRAY_SECTION_STRIDED
- && cp_lexer_next_token_is (parser->lexer, CPP_COLON))
+ && (cp_lexer_next_token_is (parser->lexer, CPP_COLON)
+ || double_scope))
{
cp_lexer_consume_token (parser->lexer);
/* We could check for C++-23 multidimensional/comma-separated
@@ -9531,7 +9563,14 @@ post_colon_parsing:
cp_parser_skip_to_closing_square_bracket (parser);
return error_mark_node;
}
- else if (!close_splice)
+ else if (parser->omp_array_section_kind == OMP_ARRAY_SECTION_STRIDED
+ && cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_SPLICE))
+ cp_lexer_consume_token (parser->lexer);
+ else if (!close_splice
+ && (parser->omp_array_section_kind
+ == OMP_ARRAY_SECTION_UNSTRIDED
+ || cp_lexer_next_token_is_not (parser->lexer,
+ CPP_CLOSE_SPLICE)))
cp_parser_require (parser, CPP_CLOSE_SQUARE, RT_CLOSE_SQUARE);
return grok_omp_array_section (input_location, postfix_expression, index,
@@ -42017,23 +42056,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
decl = error_mark_node;
}
else
- {
- /* We have a pointer DECL whose target should be
- interpreted as an array with particular dimensions,
- not "the pointer itself". So, add an indirection
- here. */
- if (type_dependent_expression_p (decl))
- decl = build_min_nt_loc (loc, INDIRECT_REF, decl);
- else
- {
- /* We're interested in the reference target. */
- decl = convert_from_reference (decl);
- decl = cp_build_fold_indirect_ref (decl);
- }
- decl
- = cp_build_omp_arrayshape_cast (loc, reshaped_to, decl,
- tf_warning_or_error);
- }
+ decl = cp_build_omp_arrayshape_cast (loc, reshaped_to, decl,
+ tf_warning_or_error);
}
/* Bare references have their own special handling, so remove
the explicit dereference added by convert_from_reference. */
@@ -42044,7 +42068,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
decl = grok_omp_array_section (loc, decl, dims[i].low_bound,
dims[i].length, dims[i].stride);
}
- else if (TREE_CODE (decl) == INDIRECT_REF)
+ else if (TREE_CODE (decl) == INDIRECT_REF && !reshaped_to)
{
bool ref_p = REFERENCE_REF_P (decl);
@@ -42073,8 +42097,16 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
else if (reshaped_to)
{
/* We're copying the whole of a reshaped array, originally a
- base pointer. Rewrite as an array section. */
+ base pointer. DECL may already be dereferenced here, so
+ reinterpret it as a flat 1D array of the total element count.
+ */
tree elems = array_type_nelts_total (reshaped_to);
+ tree elt_type = reshaped_to;
+ while (TREE_CODE (elt_type) == ARRAY_TYPE)
+ elt_type = TREE_TYPE (elt_type);
+ tree flat_type = build_array_type_nelts (elt_type, elems);
+ decl = cp_build_omp_arrayshape_cast (loc, flat_type, decl,
+ tf_warning_or_error);
decl = grok_omp_array_section (loc, decl, size_zero_node, elems,
NULL_TREE);
}
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index df377413857..5df9d53094c 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -6317,6 +6317,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
+ if (stride != NULL_TREE && TREE_CODE (stride) == INTEGER_CST
+ && (tree_int_cst_sgn (stride) == -1
+ || tree_int_cst_sgn (stride) == 0))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "negative or zero stride in array section in %qs clause",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ return error_mark_node;
+ }
if (TYPE_DOMAIN (type)
&& TYPE_MAX_VALUE (TYPE_DOMAIN (type))
&& TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (type)))
@@ -15223,6 +15232,21 @@ cp_build_omp_arrayshape_cast (location_t loc, tree type, tree arg,
return error_mark_node;
}
+ if (TREE_TYPE (arg) != NULL_TREE && POINTER_TYPE_P (TREE_TYPE (arg)))
+ {
+ /* We have a pointer DECL whose target should be interpreted as an
+ array with particular dimensions, not "the pointer itself". So,
+ add an indirection here. */
+ if (type_dependent_expression_p (arg))
+ arg = build_min_nt_loc (loc, INDIRECT_REF, arg);
+ else
+ {
+ /* We're interested in the reference target. */
+ arg = convert_from_reference (arg);
+ arg = cp_build_fold_indirect_ref (arg);
+ }
+ }
+
/* A pointer to multi-dimensional array conversion isn't normally
allowed, but we force it here for array shape operators by creating
the node directly. We also want to avoid any overloaded conversions
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c
index 05fb788e817..8f18d6b0ff7 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c
@@ -9,5 +9,5 @@
void f (int **x)
{
- #pragma omp target update to (iterator(i=0:DIM1): x[i][0:DIM2:2]) /* { dg-message "sorry, unimplemented: 'iterator' modifier combined with a noncontiguous map expression" } */
+ #pragma omp target update to (iterator(i=0:DIM1): x[i][0:DIM2:2]) /* { dg-message "sorry, unimplemented: 'iterator' modifier combined with a non-contiguous array section" } */
}
diff --git a/gcc/testsuite/g++.dg/gomp/array-section-7.C b/gcc/testsuite/g++.dg/gomp/array-section-7.C
new file mode 100644
index 00000000000..5dc87030c08
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/array-section-7.C
@@ -0,0 +1,58 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-original" }
+
+// Check parsing of double colon (::). In case of ambiguity between the scope
+// resolution operator and the array section syntax, spaces are required.
+
+int lb2 = 4;
+int len2 = 16;
+int s2 = 2;
+
+void g() {
+ int arr[128];
+ int lb = 4, len = 16, s = 2;
+
+#pragma omp target update to(arr[lb::s])
+// { dg-error "34: 'lb' is not a class" "" { target *-*-* } .-1 }
+// { dg-error "must contain at least one" "" { target *-*-* } .-2 }
+
+#pragma omp target update to(arr[::s])
+// { dg-error "36: '::s' has not been declared" "" { target *-*-* } .-1 }
+// { dg-error "must contain at least one" "" { target *-*-* } .-2 }
+
+#pragma omp target update to(arr[lb::])
+// { dg-error "34: 'lb' is not a class" "" { target *-*-* } .-1 }
+// { dg-error "38: expected unqualified-id before '\\\]' token" "" { target *-*-* } .-2 }
+// { dg-error "must contain at least one" "" { target *-*-* } .-3 }
+
+ // Valid: default values.
+#pragma omp target update to(arr[::])
+// { dg-final { scan-tree-dump "to\\(arr\\\[0\\\] \\\[len: 512\\\]\\)" "original" } }
+
+ // Valid: len2 is looked up in the global namespace and forms the length of the
+ // array section.
+#pragma omp target update to(arr[:::len2])
+#pragma omp target update to(arr[:::len2:])
+// { dg-final { scan-tree-dump-times "to\\(arr\\\[0\\\] \\\[len: \\(sizetype\\) len2 \\* 4\\\]\\)" 2 "original" } }
+
+ // Valid: s2 is looked up in the global namespace and forms the stride of the
+ // array section. The length is omitted, so it defaults to
+ // ceil((128 - 0) / s2) elements.
+#pragma omp target update to(arr[::::s2])
+// { dg-final { scan-tree-dump "map\\(to_grid:arr \\\[len: 4\\\]\\) map\\(grid_dim:0 \\\[len: 128 /\\\[cl\\\] \\(sizetype\\) s2\\\]\\) map\\(grid_stride:s2\\)" "original" } }
+
+ // Valid: lb2 is looked up in the global namespace and forms the lower bound
+ // of the array section.
+#pragma omp target update to(arr[::lb2:])
+// { dg-final { scan-tree-dump "to\\(arr\\\[SAVE_EXPR <lb2>\\\] \\\[len: \\(128 - \\(sizetype\\) SAVE_EXPR <lb2>\\) \\* 4\\\]\\)" "original" } }
+
+#pragma omp target update to(arr[::lb2::])
+// { dg-error "36: '::lb2' is not a class" "" { target *-*-* } .-1 }
+// { dg-error "41: expected unqualified-id before '\\\]' token" "" { target *-*-* } .-2 }
+// { dg-error "must contain at least one" "" { target *-*-* } .-3 }
+
+#pragma omp target update to(arr[::lb2::s2])
+// { dg-error "36: '::lb2' is not a class" "" { target *-*-* } .-1 }
+// { dg-error "must contain at least one" "" { target *-*-* } .-2 }
+
+}
diff --git a/gcc/testsuite/g++.dg/gomp/array-section-8.C b/gcc/testsuite/g++.dg/gomp/array-section-8.C
new file mode 100644
index 00000000000..622e100be29
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/array-section-8.C
@@ -0,0 +1,36 @@
+// { dg-do compile { target c++26 } }
+// { dg-additional-options "-freflection -fdump-tree-original" }
+
+// Check that the OpenMP stride array-section notation
+// '[ lower : length : stride ]' is not confused with C++26's splice specifier,
+// including when the latter is used to supply one of the section's components.
+
+constexpr int lb0 = 4;
+constexpr auto rlb = ^^lb0;
+constexpr int s0 = 2;
+constexpr auto rs = ^^s0;
+
+void f() {
+ int arr[128];
+ int lb = 4, len = 16, s = 2;
+
+ // baseline: all three components explicit, nothing splice-shaped.
+#pragma omp target update to(arr[lb:len:s])
+
+ // splice-expression supplying the lower bound.
+#pragma omp target update to(arr[[:rlb:]:len:s])
+
+ // splice-expression supplying the stride.
+#pragma omp target update to(arr[lb:len:[:rs:]])
+
+ // Stride left empty: the trailing ':' + ']' is lexed as a single
+ // CPP_CLOSE_SPLICE token.
+#pragma omp target update to(arr[lb:len:])
+#pragma omp target update to(arr[:len:])
+}
+
+// { dg-final { scan-tree-dump "map\\(to_grid:arr \\\[len: 4\\\]\\) map\\(grid_dim:SAVE_EXPR <lb> \\\[len: len\\\]\\) map\\(grid_stride:s\\)" "original" } }
+// { dg-final { scan-tree-dump "map\\(to_grid:arr \\\[len: 4\\\]\\) map\\(grid_dim:4 \\\[len: len\\\]\\) map\\(grid_stride:s\\)" "original" } }
+// { dg-final { scan-tree-dump "map\\(to_grid:arr \\\[len: 4\\\]\\) map\\(grid_dim:SAVE_EXPR <lb> \\\[len: len\\\]\\) map\\(grid_stride:2\\)" "original" } }
+// { dg-final { scan-tree-dump "to\\(arr\\\[SAVE_EXPR <lb>\\\] \\\[len: \\(sizetype\\) len \\* 4\\\]\\)" "original" } }
+// { dg-final { scan-tree-dump "to\\(arr\\\[0\\\] \\\[len: \\(sizetype\\) len \\* 4\\\]\\)" "original" } }
diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-section-12.C b/gcc/testsuite/g++.dg/gomp/bad-array-section-12.C
new file mode 100644
index 00000000000..1a5dd470076
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/bad-array-section-12.C
@@ -0,0 +1,32 @@
+// { dg-do compile }
+
+/* A stride specifier is only valid for "to"/"from" clauses (including on
+ "target update"), not for "map" clauses -- not even a "map" clause
+ using the "to"/"from" modifier. */
+
+int main ()
+{
+ int x[10];
+
+#pragma omp target enter data map(to: x)
+
+#pragma omp target map(x[0:5:2])
+ // { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 }
+ // { dg-error {expected '\)' before ':' token} "" { target *-*-* } .-2 }
+ // { dg-error "expected an OpenMP clause before '\\\]' token" "" { target *-*-* } .-3 }
+ ;
+
+#pragma omp target enter data map(to: x[0:5:2])
+ // { dg-error {expected '\]' before ':' token} "" { target *-*-* } .-1 }
+ // { dg-error {expected '\)' before ':' token} "" { target *-*-* } .-2 }
+ // { dg-error "expected an OpenMP clause before '\\\]' token" "" { target *-*-* } .-3 }
+
+ /* These are fine: "to"/"from" clauses on "target update" do accept a
+ stride. */
+#pragma omp target update to(x[0:5:2])
+#pragma omp target update from(x[0:5:2])
+
+#pragma omp target exit data map(release: x)
+
+ return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-section-13.C b/gcc/testsuite/g++.dg/gomp/bad-array-section-13.C
new file mode 100644
index 00000000000..f8640907f0e
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/bad-array-section-13.C
@@ -0,0 +1,16 @@
+// { dg-do compile }
+
+void f()
+{
+ int x[10];
+ int y[10][10][10];
+ #pragma omp target update to(x[0:0:-1])
+ // { dg-error "negative or zero stride in array section in 'to' clause" "" { target *-*-* } .-1 }
+ // { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 }
+ #pragma omp target update to(x[0:0:0])
+ // { dg-error "negative or zero stride in array section in 'to' clause" "" { target *-*-* } .-1 }
+ // { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 }
+ #pragma omp target update to(y[0:0:0])
+ // { dg-error "negative or zero stride in array section in 'to' clause" "" { target *-*-* } .-1 }
+ // { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 }
+}
diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C
index 02d7de6088e..4b077ca86a9 100644
--- a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C
+++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C
@@ -19,7 +19,7 @@ void foo ()
// { dg-error {high bound '101' above array section size in 'from' clause} "" { target *-*-* } .-1 }
#pragma omp target update to(([10][10]) ptr[0:9:-1][0:9])
-// { dg-error {length '9' with stride '-1' above array section size in 'to' clause} "" { target *-*-* } .-1 }
+// { dg-error {negative or zero stride in array section in 'to' clause} "" { target *-*-* } .-1 }
}
int main()
@@ -41,7 +41,7 @@ int main()
// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 }
#pragma omp target update to(([10][10]) ptr[0:9:-1][0:9])
-// { dg-error {length '9' with stride '-1' above array section size in 'to' clause} "" { target *-*-* } .-1 }
+// { dg-error {negative or zero stride in array section in 'to' clause} "" { target *-*-* } .-1 }
// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 }
foo<char> ();
diff --git a/gcc/tree.def b/gcc/tree.def
index 0e77005abbd..362295fdd70 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1422,7 +1422,11 @@ DEFTREECODE (OMP_ATOMIC_CAPTURE_NEW, "omp_atomic_capture_new", tcc_statement, 2)
/* OpenMP clauses. */
DEFTREECODE (OMP_CLAUSE, "omp_clause", tcc_exceptional, 0)
-/* An OpenMP array section. */
+/* An OpenMP extended array section.
+ Operand 0: The base array declaration.
+ Operand 1: The lower bound.
+ Operand 2: The length.
+ Operand 3: The stride. */
DEFTREECODE (OMP_ARRAY_SECTION, "omp_array_section", tcc_expression, 4)
/* OpenMP variant construct selector, used only in the middle end in the
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-14.C b/libgomp/testsuite/libgomp.c++/array-shaping-14.C
new file mode 100644
index 00000000000..606464aa1f8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-14.C
@@ -0,0 +1,65 @@
+// { dg-do run { target offload_device_nonshared_as } }
+
+/* Correctness test for:
+ - omitted low_bound/length/stride defaulting in an OpenMP array
+ section (low_bound -> 0, length -> ceil((size-low_bound)/stride),
+ stride -> 1), and
+ - the array-shaping operator applied directly to a bare pointer
+ (([10][8][6])p), which must dereference p before reinterpreting
+ its type.
+ */
+
+int main() {
+ int x[10][8][6];
+ int *p = (int *) __builtin_malloc (sizeof (int) * 10 * 8 * 6);
+
+ for (int i = 0; i < 10; i++)
+ for (int j = 0; j < 8; j++)
+ for (int k = 0; k < 6; k++)
+ {
+ x[i][j][k] = 100*i + 10*j + k;
+ p[i*8*6 + j*6 + k] = 100*i + 10*j + k;
+ }
+
+ #pragma omp target enter data map(to: x, p[:10*8*6])
+
+ #pragma omp target map(present, alloc: x, p[:10*8*6])
+ for (int i = 0; i < 10; i++)
+ for (int j = 0; j < 8; j++)
+ for (int k = 0; k < 6; k++)
+ {
+ x[i][j][k] *= 23;
+ p[i*8*6 + j*6 + k] *= 23;
+ }
+
+ /* Single update, covering both x and the array-shaping-cast p, each
+ with all bounds/length omitted and only a stride given. */
+ #pragma omp target update from(x[: :4][: :3][: :2], \
+ (([10][8][6])p)[: :4][: :3][: :2])
+
+ for (int i = 0; i < 10; i++)
+ for (int j = 0; j < 8; j++)
+ for (int k = 0; k < 6; k++)
+ {
+ bool dev_value = (i % 4 == 0) && (j % 3 == 0) && (k % 2 == 0);
+ int factor = dev_value ? 23 : 1;
+ if (x[i][j][k] != factor * (100*i + 10*j + k))
+ __builtin_abort ();
+ if (p[i*8*6 + j*6 + k] != factor * (100*i + 10*j + k))
+ __builtin_abort ();
+ }
+
+ #pragma omp target exit data map(from: x, p[:10*8*6])
+
+ for (int i = 0; i < 10; i++)
+ for (int j = 0; j < 8; j++)
+ for (int k = 0; k < 6; k++)
+ {
+ if (x[i][j][k] != 23 * (100*i + 10*j + k))
+ __builtin_abort ();
+ if (p[i*8*6 + j*6 + k] != 23 * (100*i + 10*j + k))
+ __builtin_abort ();
+ }
+
+ __builtin_free (p);
+}
--
2.55.0
