Hi Tobias,
Thanks for the review. See my replies inline below and the attached
fixup patch.
I'll post the C++-specific fixes in a separate patch under the relevant
sub-thread.
On 24/08/2026 16:29, Tobias Burnus wrote:
Hi PA,
Paul-Antoine Arras wrote:
From: Julian Brown<[email protected]>
Following the similar support for C++, here is the C implementation for
the OpenMP 5.0 array-shaping operator, and for strided and rectangular
updates for "target update".
Doing the FE review in reverse, starting with the C FE review ...
I guess some comments also apply to the C++ FE.
* * *
Zeroth: Thanks for taking care of upstreaming Julian's patch and
doing some fixes and cleanup there!
First, can you check
https://github.com/OpenMP/Examples/blob/main/devices/sources/
array_shaping.1.c
unless I screwed up, it fails here with the C compiler but not the C++
compiler:
devices/sources/array_shaping.1.c: In function ‘array_shaping’:
devices/sources/array_shaping.1.c:26:15: internal compiler error: in
c_set_type_bits, at c/c-typeck.cc:466
26 | #pragma omp target update from( (([nx][ny+2])a)[0:nx:2]
[1], \
| ^~~
0x2a6d69f internal_error(char const*, ...)
../../../repos/gcc/gcc/diagnostic-global-context.cc:787
0xa761f9 fancy_abort(char const*, int, char const*)
../../../repos/gcc/gcc/diagnostics/context.cc:1813
0x7bd7f8 c_set_type_bits
../../../repos/gcc/gcc/c/c-typeck.cc:466
0xab82a5 build_omp_array_section(unsigned long, tree_node*, tree_node*,
tree_node*, tree_node*)
../../../repos/gcc/gcc/c/c-typeck.cc:3776
Good catch! That was a copy-paste mistake from C++ to C: in
create_omp_arrayshape_type, build_array_type was called instead of
*c_*build_array_type.
Added testcase gcc/testsuite/gcc.dg/gomp/array-shaping-1.c.
* * *
Talking about ICE, the OpenMP_VV Fortran tests, tests/5.1/target_update/
test_target_update_iterator.F90
https://github.com/OpenMP-Validation-and-Verification/OpenMP_VV/blob/
master/tests/5.1/target_update/test_target_update_iterator.F90
fails once the Fortran patch is applied - it works without the patch and
does not
use the new features. Namely:
internal compiler error: in gfc_conv_is_contiguous_expr, at fortran/
trans-intrinsic.cc:2334
0x2b0555f internal_error(char const*, ...)
../../../repos/gcc/gcc/diagnostic-global-context.cc:787
0xa6fe0b fancy_abort(char const*, int, char const*)
../../../repos/gcc/gcc/diagnostics/context.cc:1813
0x7e0eeb gfc_conv_is_contiguous_expr(gfc_se*, gfc_expr*)
../../../repos/gcc/gcc/fortran/trans-intrinsic.cc:2334
0xc609ac gfc_trans_omp_clauses
../../../repos/gcc/gcc/fortran/trans-openmp.cc:5911
0xc7126d gfc_trans_omp_target_update
../../../repos/gcc/gcc/fortran/trans-openmp.cc:10259
Noted. I'll handle that as part of the Fortran patch sub-thread.
* * *
The following seems to be completely valid C:
void f() {
int A[10], B[10][10];
#pragma omp target update to(A[::], B[::][::1])
}
but it gets rejected with:
r.c:3:32: error: expected expression before ‘::’ token
3 | #pragma omp target update to(A[::], B[::][::1])
IMHO this should just work in C.
Agreed. The double colon is parsed as CPP_SCOPE, which wasn't handled.
Now fixed and added testcase:
gcc/testsuite/c-c++-common/gomp/array-section-1.c.
* * *
For C++, I think 'A[::]' is supposed to be likewise valid
but gets rejected:
r.c:4:34: error: expected id-expression before ‘]’ token
4 | #pragma omp target update to(A[::])
| ^
Same root cause: CPP_SCOPE was mishandled. This is resolved by the fixup
patch I posted in the C++ sub-thread.
Talking about C++, the following works:
int n;
void f() {
int A[10], B[10][10];
#pragma omp target update to(A[::n])
}
as expected, but failing is:
r.c:4:34: error: expected id-expression before ‘]’ token
4 | #pragma omp target update to(A[::], B[0][: : ])
| ^
r.c:4:44: error: expected primary-expression before ‘:’ token
4 | #pragma omp target update to(A[::], B[0][: : ])
| ^
I think this should work by the same reasoning as for '[: :]'.
Ditto for C++. Now fixed for C and added testcase.
And I am pretty sure that '[::n :]' and '[:]' works with map
clauses with g++.
* * *
Also failing - but not really different:
r.c:4:49: error: expected primary-expression before ‘:’ token
4 | #pragma omp target update to(A[1:2:], B[:1:][1: : ])
I think some of them - esp. the last one - I already reported.
Ditto.
* * *
And failing but trickier due to the sequence of > 3 colons:
[Cf. OpenMP 6.x and discussions, cf. OpenMP spec issue #4740]:
r.c:4:34: error: expected id-expression before ‘:’ token
4 | #pragma omp target update to(A[:::n], B[:::n][: :: n : ])
| ^
r.c:4:43: error: expected id-expression before ‘:’ token
4 | #pragma omp target update to(A[:::n], B[:::n][: :: n : ])
| ^
Ditto for C++. Invalid in C.
where the A[::n] is an array element (this works) and
:::n is the length '::n' and ::::n is the stride '::n'.
* * *
2026-08-13 Paul-Antoine Arras<[email protected]>
gcc/c/
* c-tree.h (c_omp_array_section_p): Replace bool with...
(c_omp_array_section_kind): ...this 3-state unsigned char.
(OMP_ARRAY_SECTION_NONE, OMP_ARRAY_SECTION_UNSTRIDED,
OMP_ARRAY_SECTION_STRIDED): Define.
* c-parser.cc (c_parser_postfix_expression_after_primary): Only
accept a 2rd colon (stride) for to/from clauses, not map.
(c_parser_omp_variable_list): Only count OMP_ARRAY_SECTION
layers, not every ARRAY_REF layer, in a to/from clause's base
decl.
* c-typeck.cc (handle_omp_array_sections): Handle a genuinely
discontiguous update at any access-chain depth.
(c_finish_omp_clauses): Splice the GOMP_MAP_TO_GRID/FROM_GRID
replacement clause in at *PC when DISCONTIGUOUS == 2.
I think it would be good to have a testcase.
Added testcase gcc/testsuite/c-c++-common/gomp/array-section-2.c.
* * *
Additionally, I wonder whether it wouldn't be cleaner to have an
'enum ... : unsigned char' instead of a bare 'unsigned char' – it
adds some type safety and also makes it easier to see what it is about.
The patch does in c-tree.h:
-extern bool c_omp_array_section_p;
+#define OMP_ARRAY_SECTION_NONE 0
+#define OMP_ARRAY_SECTION_UNSTRIDED 1
+#define OMP_ARRAY_SECTION_STRIDED 2
+extern unsigned char c_omp_array_section_kind;
Updated as suggested:
enum c_omp_array_section_kind : unsigned char {
OMP_ARRAY_SECTION_NONE = 0,
OMP_ARRAY_SECTION_UNSTRIDED = 1,
OMP_ARRAY_SECTION_STRIDED = 2
};
extern enum c_omp_array_section_kind c_omp_array_section_kind;
* * *
Can you update the following comment in c-typeck.c?
'True' no longer makes sense:
/* True when parsing OpenMP map clause. */
-bool c_omp_array_section_p;
+unsigned char c_omp_array_section_kind;
Code wise, there is is still the assumption that the following
condition works:
c-typeck.cc: if (TREE_TYPE (ref) == error_mark_node && !
c_omp_array_section_kind)
i.e. that '(bool) OMP_ARRAY_SECTION_NONE' evaluates to false.
I think it would be cleaner to compare against NONE, esp. if/once
changing to an enum.
Updated comment and removed implicit conversions of
c_omp_array_section_kind.
* * *
+create_omp_arrayshape_type (tree expr, vec<tree> *omp_shape_dims)
+{
...
+ error ("OpenMP array shaping operator with non-pointer argument");
Can we have an 'error_at' here? I guess "EXPR_LOCATION (expr)" will work.
[This also applies to the C++ FE code.]
In most cases, expr is a VAR_DECL so EXPR_LOCATION won't work. Instead I
have now passed cast_loc into create_omp_arrayshape_type.
Also added test: gcc/testsuite/c-c++-common/gomp/bad-array-shaping-1.c.
* * *
@@ -15846,7 +15892,7 @@ c_finish_omp_cancellation_point (location_t
loc, tree clauses)
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
bool &maybe_zero_len, unsigned int &first_non_one,
- enum c_omp_region_type ort)
+ enum c_omp_region_type ort, int *discontiguous)
{
...
+/* Handle array sections for clause C. On entry *DISCONTIGUOUS is 0
if array
+ section must be contiguous, 1 if it can be discontiguous, and in
the latter
+ case it is set to 2 on exit if it is determined to be
discontiguous during
+ the function's execution. */
static bool
-handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
+handle_omp_array_sections (tree &c, enum c_omp_region_type ort,
+ int *discontiguous = NULL)
First and unrelated to this patch, I think it would be helpful to document
the return value of these two functions (handle_omp_array_sections_1 and
handle_omp_array_sections). Namely,
The aux _1 function returns a tree - the other a Boolean.
For the latter 'true' means an error, the former returns
error_mark_node in case of an error, for non OpenMP array sections
essentially 't' (except for a convert_lvalue_to_rvalue conversion
in a corner case) - and for array sections either build_array_ref or
build_omp_array_section.
Updated both comments to mention return values.
* * *
I think the aux function needs also comment documenting the discontiguous
argument.
For the handle_omp_array_sections_1 aux function:
If known to be discontiguous:
* Print an error if discontiguous is a nullptr or *discontiguous == 0
* Otherwise, set *discontiguous = 2
If discontiguous && *discontiguous != 0 - build_omp_array_section.
For handle_omp_array_sections:
If discontiguous && *discontiguous
* When the array might be noncontigous, set *discontiguous = 2
* If at the end *discontiguous remains with value != 2, i.e. the
array section is known to be contiguous, convert the array section
back to an build_array_ref.
Hence:
For the aux function, '2' strictly means that it is known to
be noncontiguous.
For the other function, '2' means that it might be noncontigous,
'1' that it is known to be noncontigous, and '0' that semantic
requires noncontigous (with a compile time check, but still could
go wrong at runtime).
Noncontiguous can be either if stride != 1 [with some exceptions] but
also for multidimension variables for, e.g. 'arr[1:n][0:3]'
if 'int arr[...][size]' and size > 3 [except n == 1].
Updated doc comments for discontiguous in both functions.
* * *
I wonder whether it wouldn't be cleaner to make DISCONTIGUOUS
required in the aux function. Using the current argument wording
for handle_omp_array_sections could be reused.
Agreed.
Otherwise, it should mention that nullptr and *discontigous == 0
both mean the same. - The nullptr bit could be removed, if the
pointer argument is alwas passed (like in the variant below):
* * *
Regarding the required part, that's because I wonder whether it is
cleaner to use ...
... in handle_omp_array_sections, the argument 'bool *discontiguous_p'
with:
int discontigous = discontiguous_p && *discontiguous_p ? 1 : 0;
...
// Discontiguous permitted but known to be contiguous.
if (discontigous == 1)
discontigous_p = false;
and passing this 'discontiguous' as argument to the aux function (i.e
it will never be a nullptr in the aux function).
Yes, probably cleaner. Let's do it this way.
Alternatively, if keeping it as integer, I wonder whether it would
be cleaner to set 'discontiguous' to '0' if known to be known to be
contiguous instead of keeping the old value.
If keeping it as integer: It should be made clear that '2' implies
that it might be discontigous (but it might be only known at runtime)
while the original value (or if set to zero, the value 0) means that
it is known to be contiguous.
And in in either case (bool* or int*), it should be made clear that
a nullptr acts like setting the value to 0 / false.
Mentioned in the top comment.
* * *
If we know that *discontiguous can only be 0 or two, the following
+ if (discontiguous && *discontiguous != 2)
+ first = omp_array_section_low_bound (OMP_CLAUSE_LOCATION (c),
first);
could be replaced by '*discontigous == 1' - as 0 won't create an array
section
and for 2 we want to retain it.
Done.
* * *
Actually, I think my claim that _1 only sets it to 2 when it is known to
be noncontiguous is not quite right:
/* If there is a pointer type anywhere but in the very first
array-section-subscript, the array section could be non-
contiguous. */
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
&& TREE_CODE (TREE_OPERAND (t, 0)) == OMP_ARRAY_SECTION)
{
/* If any prior dimension has a non-one length, then deem this
array section as non-contiguous. */
for (tree d = TREE_OPERAND (t, 0);
TREE_CODE (d) == OMP_ARRAY_SECTION;
d = TREE_OPERAND (d, 0))
{
tree d_length = TREE_OPERAND (d, 2);
- if (d_length == NULL_TREE || !integer_onep (d_length))
+ tree d_stride = TREE_OPERAND (d, 3);
+ if (d_length == NULL_TREE || !integer_onep (d_length)
+ || (d_stride && !integer_onep (d_stride)))
Assume:
to(arr[:1:m]) or to(arr[:1:2])
Those are the elements:
{ lower-bound, ... , lower-bound + ((length - 1) * stride) }
Namely, only '{ lower-bound }' remains - such that the array section
can still be contiguous.
On the other hand, for:
to(arr[:2]) or to(arr[:n])
the condition is already true - independent whether there is any stride
or not.
(for n == 1 it would be still contiguous - but that's not handled in
this very
special case.)
Thus, I think we can remove the 'd_stride' here - as it either leads
either to
a false positive or is redundant.
Indeed.
* * *
Hence, I think for
+handle_omp_array_sections (tree &c, enum c_omp_region_type ort,
....
+ if (!integer_onep (stride)
+ || (higher_discontiguous
+ && (!integer_zerop (low_bound)
+ || !full_span)))
+ *discontiguous = 2;
+
+ if (!integer_onep (stride)
+ || !integer_zerop (low_bound)
+ || !full_span)
+ higher_discontiguous = true;
IMHO, it should be for both:
if ((!integer_onep (stride) && !integer_onep (length)
for the same reasons (if length is one, the stride >= 1
does not matter).
Adjusted both conditions.
* * *
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ && !(discontiguous && *discontiguous == 2))
return false;
I wonder whether a comment would help:
// Done, except for MAP clauses and for array sections that may be
noncontiguous
Added comment.
At least I find it otherwise a bit hard to read with the ! (not) and
more complex
conditions - especially as at a glance, there is no real relation
between one and
the other. (Namely: Why aren't TO/FROM handled the same way as map, if
contiguous?
Answer: See longer comment below.)
Side note: c_omp_address_inspector (which is later called) actually
turns the
TO/FROM clause to a MAP clause of map kind GOMP_MAP_{TO,FROM}_GRID;
however,
that happens after this check (and is actually the reason for continuing).
* * *
Back to the aux function:
@@ -16198,14 +16277,42 @@ handle_omp_array_sections_1 (tree c, tree t,
vec<tree> &types,
- ret = build_array_ref (OMP_CLAUSE_LOCATION (c), ret, low_bound);
+ /* NOTE: Stride/length are discarded for affinity/depend here. */
+ if (discontiguous
+ && *discontiguous
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND)
+ ret = build_omp_array_section (OMP_CLAUSE_LOCATION (c), ret,
low_bound,
+ length, stride);
+ else
+ ret = build_array_ref (OMP_CLAUSE_LOCATION (c), ret, low_bound);
return ret;
Can we remove the comment and clause-code checks here? Namely:
I have to admit that I find the comment and the check for
affinity/depend more confusing than helpful.
All callers either pass no discontiguous argument (→ NULL default),
except for one:
+ int discontiguous
+ = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM);
+ if (handle_omp_array_sections (c, ort, &discontiguous))
And this one explicitly handles those two clauses.
The reason that I find it odd is that the aux function is also called
for a much of other clauses, e.f., OMP_CLAUSE_MAP or OMP_CLAUSE__CACHE_
for which the same applies.
Good catch! Removed the extra predicates.
Thanks,
--
PA
From a4bd15474ba367d9d6b6e3db92b076ad664f4af7 Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <[email protected]>
Date: Thu, 27 Aug 2026 17:49:59 +0200
Subject: [PATCH] OpenMP: Fix strided and shaped-array update issues for C
This is a fixup commit for 'OpenMP: Array shaping operator and strided "target
update" for C', addressing issues found during review of that patch.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_braced_init): Adjust for
c_omp_array_section_kind becoming a fixed-underlying-type enum.
(c_parser_conditional_expression): Likewise.
(c_parser_cast_expression): Likewise, and pass a location to
create_omp_arrayshape_type.
(c_parser_postfix_expression): Adjust for
c_omp_array_section_kind becoming a fixed-underlying-type enum.
(c_parser_postfix_expression_after_primary): Also recognise the
"::" token as a stride-only array-section separator when parsing
OMP_ARRAY_SECTION_STRIDED sections.
(c_parser_expr_list): Adjust for c_omp_array_section_kind becoming
a fixed-underlying-type enum.
(c_parser_omp_variable_list): Likewise.
* c-tree.h (OMP_ARRAY_SECTION_NONE, OMP_ARRAY_SECTION_UNSTRIDED,
OMP_ARRAY_SECTION_STRIDED, enum c_omp_array_section_kind): Turn the
array-section-kind macros and the unsigned char global into a
fixed-underlying-type enum.
(create_omp_arrayshape_type): Add a location_t parameter.
* c-typeck.cc (c_omp_array_section_kind): Adjust for the new enum
type.
(create_omp_arrayshape_type): Take a location_t parameter and use
it for diagnostics; build the result type with c_build_array_type
instead of build_array_type.
(build_external_ref): Adjust for the new enum type.
(handle_omp_array_sections_1): Document the DISCONTIGUOUS parameter
and return value; drop the redundant stride check when detecting a
non-contiguous pointer-typed dimension; simplify now that
DISCONTIGUOUS is always non-null.
(handle_omp_array_sections): Change the DISCONTIGUOUS parameter
from int * to bool *; only treat a non-unit stride as
discontiguous when the length is also not one; add a comment to
the early-return check.
(c_finish_omp_clauses): Change the local discontiguous variable to
bool.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/array-section-1.c: New test.
* c-c++-common/gomp/array-section-2.c: New test.
* c-c++-common/gomp/bad-array-shaping-1.c: New test.
* gcc.dg/gomp/array-shaping-1.c: New test.
---
gcc/c/c-parser.cc | 158 +++++++++++-------
gcc/c/c-tree.h | 12 +-
gcc/c/c-typeck.cc | 87 ++++++----
.../c-c++-common/gomp/array-section-1.c | 20 +++
.../c-c++-common/gomp/array-section-2.c | 32 ++++
.../c-c++-common/gomp/bad-array-shaping-1.c | 12 ++
gcc/testsuite/gcc.dg/gomp/array-shaping-1.c | 17 ++
7 files changed, 235 insertions(+), 103 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-1.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-2.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/bad-array-shaping-1.c
create mode 100644 gcc/testsuite/gcc.dg/gomp/array-shaping-1.c
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index f8a6a98118b..512ad47bfb1 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -6636,9 +6636,10 @@ c_parser_braced_init (c_parser *parser, tree type, bool nested_p,
location_t brace_loc = c_parser_peek_token (parser)->location;
gcc_obstack_init (&braced_init_obstack);
gcc_assert (c_parser_next_token_is (parser, CPP_OPEN_BRACE));
- unsigned char save_c_omp_array_section_kind = c_omp_array_section_kind;
+ enum c_omp_array_section_kind save_c_omp_array_section_kind
+ = c_omp_array_section_kind;
bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
- c_omp_array_section_kind = false;
+ c_omp_array_section_kind = OMP_ARRAY_SECTION_NONE;
c_omp_array_shaping_op_p = false;
bool zero_init_padding_bits = false;
matching_braces braces;
@@ -10106,7 +10107,8 @@ c_parser_conditional_expression (c_parser *parser, struct c_expr *after,
{
struct c_expr cond, exp1, exp2, ret;
location_t start, cond_loc, colon_loc;
- unsigned char save_c_omp_array_section_kind = c_omp_array_section_kind;
+ enum c_omp_array_section_kind save_c_omp_array_section_kind
+ = c_omp_array_section_kind;
bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
gcc_assert (!after || c_dialect_objc ());
@@ -10115,7 +10117,7 @@ c_parser_conditional_expression (c_parser *parser, struct c_expr *after,
if (c_parser_next_token_is_not (parser, CPP_QUERY))
return cond;
- c_omp_array_section_kind = false;
+ c_omp_array_section_kind = OMP_ARRAY_SECTION_NONE;
c_omp_array_shaping_op_p = false;
if (cond.value != error_mark_node)
start = cond.get_start ();
@@ -10619,9 +10621,10 @@ c_parser_cast_expression (c_parser *parser, struct c_expr *after)
if (c_parser_next_token_is (parser, CPP_OPEN_PAREN)
&& c_token_starts_compound_literal (c_parser_peek_2nd_token (parser)))
{
- unsigned char save_c_omp_array_section_kind = c_omp_array_section_kind;
+ enum c_omp_array_section_kind save_c_omp_array_section_kind
+ = c_omp_array_section_kind;
bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
- c_omp_array_section_kind = false;
+ c_omp_array_section_kind = OMP_ARRAY_SECTION_NONE;
c_omp_array_shaping_op_p = false;
struct c_declspecs *scspecs;
struct c_type_name *type_name;
@@ -10679,9 +10682,10 @@ c_parser_cast_expression (c_parser *parser, struct c_expr *after)
&& c_parser_next_token_is (parser, CPP_OPEN_PAREN)
&& c_parser_peek_2nd_token (parser)->type == CPP_OPEN_SQUARE)
{
- unsigned char save_c_omp_array_section_kind = c_omp_array_section_kind;
+ enum c_omp_array_section_kind save_c_omp_array_section_kind
+ = c_omp_array_section_kind;
bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
- c_omp_array_section_kind = false;
+ c_omp_array_section_kind = OMP_ARRAY_SECTION_NONE;
c_omp_array_shaping_op_p = false;
auto_vec<tree, 4> omp_shape_dims;
struct c_expr expr, ret;
@@ -10711,7 +10715,7 @@ c_parser_cast_expression (c_parser *parser, struct c_expr *after)
expr = convert_lvalue_to_rvalue (expr_loc, expr, false, true);
}
tree arrtype
- = create_omp_arrayshape_type (expr.value, &omp_shape_dims);
+ = create_omp_arrayshape_type (cast_loc, expr.value, &omp_shape_dims);
ret.value = build1_loc (cast_loc, VIEW_CONVERT_EXPR, arrtype,
expr.value);
if (ret.value && expr.value)
@@ -11979,7 +11983,8 @@ c_parser_postfix_expression (c_parser *parser)
/* A statement expression. */
tree stmt;
location_t brace_loc;
- unsigned char save_c_omp_array_section_kind = c_omp_array_section_kind;
+ enum c_omp_array_section_kind save_c_omp_array_section_kind
+ = c_omp_array_section_kind;
bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
c_parser_consume_token (parser);
brace_loc = c_parser_peek_token (parser)->location;
@@ -11997,7 +12002,7 @@ c_parser_postfix_expression (c_parser *parser)
expr.set_error ();
break;
}
- c_omp_array_section_kind = false;
+ c_omp_array_section_kind = OMP_ARRAY_SECTION_NONE;
c_omp_array_shaping_op_p = false;
stmt = c_begin_stmt_expr ();
c_parser_compound_statement_nostart (parser);
@@ -14118,11 +14123,15 @@ warn_for_abs (location_t loc, tree fndecl, tree arg)
"of value", fndecl, atype, ftype);
}
-
/* Parse a postfix expression after the initial primary or compound
literal; that is, parse a series of postfix operators.
- EXPR_LOC is the location of the primary expression. */
+ EXPR_LOC is the location of the primary expression.
+
+ With parser->omp_array_section_kind, it also handles OpenMP
+ array sections of the type [ index : length : stride ] where index,
+ length and stride are optional. Note that an absent length might be
+ parsed as CPP_SCOPE ('::'). */
static struct c_expr
c_parser_postfix_expression_after_primary (c_parser *parser,
@@ -14147,60 +14156,81 @@ c_parser_postfix_expression_after_primary (c_parser *parser,
switch (c_parser_peek_token (parser)->type)
{
case CPP_OPEN_SQUARE:
- /* Array reference. */
- c_parser_consume_token (parser);
- idx = len = NULL_TREE;
- if (!c_omp_array_section_kind
- || c_parser_next_token_is_not (parser, CPP_COLON))
- idx = c_parser_expression (parser).value;
-
- if (c_omp_array_section_kind
- && c_parser_next_token_is (parser, CPP_COLON))
- {
- tree stride = NULL_TREE;
+ {
+ /* Array reference. */
+ c_parser_consume_token (parser);
+ idx = len = NULL_TREE;
+ if (c_omp_array_section_kind == OMP_ARRAY_SECTION_NONE
+ || (c_parser_next_token_is_not (parser, CPP_COLON)
+ && c_parser_next_token_is_not (parser, CPP_SCOPE)))
+ idx = c_parser_expression (parser).value;
+ bool double_colon
+ = c_parser_next_token_is (parser, CPP_SCOPE)
+ || (c_parser_next_token_is (parser, CPP_COLON)
+ && c_parser_peek_2nd_token (parser)->type == CPP_COLON);
+ if (double_colon && c_parser_next_token_is (parser, CPP_COLON))
c_parser_consume_token (parser);
- if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE))
- len = c_parser_expression (parser).value;
- if (c_omp_array_section_kind == OMP_ARRAY_SECTION_STRIDED
- && c_parser_next_token_is (parser, CPP_COLON))
- {
- c_parser_consume_token (parser);
- if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE))
- stride = c_parser_expression (parser).value;
- }
+ if (c_omp_array_section_kind != OMP_ARRAY_SECTION_NONE
+ && c_parser_next_token_is (parser, CPP_COLON) && !double_colon)
+ {
+ tree stride = NULL_TREE;
+
+ c_parser_consume_token (parser);
+ if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE))
+ len = c_parser_expression (parser).value;
+
+ if (c_omp_array_section_kind == OMP_ARRAY_SECTION_STRIDED
+ && c_parser_next_token_is (parser, CPP_COLON))
+ {
+ c_parser_consume_token (parser);
+ if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE))
+ stride = c_parser_expression (parser).value;
+ }
- expr.value = build_omp_array_section (op_loc, expr.value, idx,
- len, stride);
- }
- else
- {
- if (c_omp_has_array_shape_p)
- /* If we have an array-shaping operator, we may not be able to
- represent a well-formed ARRAY_REF here, because we are
- coercing the type of the innermost array base and the
- original type may not be compatible. Use the
- OMP_ARRAY_SECTION code instead. We also want to explicitly
- avoid creating INDIRECT_REFs for pointer bases, because
- that can lead to parsing ambiguities (see
- c_parser_omp_variable_list). */
expr.value = build_omp_array_section (op_loc, expr.value, idx,
- size_one_node, NULL_TREE);
- else
- expr.value = build_array_ref (op_loc, expr.value, idx);
- }
+ len, stride);
+ }
+ else if (c_omp_array_section_kind == OMP_ARRAY_SECTION_STRIDED
+ && double_colon)
+ {
+ tree stride = NULL_TREE;
+ c_parser_consume_token (parser);
+ if (c_parser_next_token_is_not (parser, CPP_CLOSE_SQUARE))
+ stride = c_parser_expression (parser).value;
+ expr.value = build_omp_array_section (op_loc, expr.value, idx,
+ len, stride);
+ }
+ else
+ {
+ if (c_omp_has_array_shape_p)
+ /* If we have an array-shaping operator, we may not be able to
+ represent a well-formed ARRAY_REF here, because we are
+ coercing the type of the innermost array base and the
+ original type may not be compatible. Use the
+ OMP_ARRAY_SECTION code instead. We also want to explicitly
+ avoid creating INDIRECT_REFs for pointer bases, because
+ that can lead to parsing ambiguities (see
+ c_parser_omp_variable_list). */
+ expr.value
+ = build_omp_array_section (op_loc, expr.value, idx,
+ size_one_node, NULL_TREE);
+ else
+ expr.value = build_array_ref (op_loc, expr.value, idx);
+ }
- c_parser_skip_until_found (parser, CPP_CLOSE_SQUARE,
- "expected %<]%>");
+ c_parser_skip_until_found (parser, CPP_CLOSE_SQUARE,
+ "expected %<]%>");
- start = expr.get_start ();
- finish = parser->tokens_buf[0].location;
- set_c_expr_source_range (&expr, start, finish);
- expr.original_code = ERROR_MARK;
- expr.original_type = NULL;
- expr.m_decimal = 0;
- break;
+ start = expr.get_start ();
+ finish = parser->tokens_buf[0].location;
+ set_c_expr_source_range (&expr, start, finish);
+ expr.original_code = ERROR_MARK;
+ expr.original_type = NULL;
+ expr.m_decimal = 0;
+ break;
+ }
case CPP_OPEN_PAREN:
/* Function call. */
{
@@ -14541,9 +14571,10 @@ c_parser_expr_list (c_parser *parser, bool convert_p, bool fold_p,
vec<tree, va_gc> *orig_types;
struct c_expr expr;
unsigned int idx = 0;
- unsigned char save_c_omp_array_section_kind = c_omp_array_section_kind;
+ enum c_omp_array_section_kind save_c_omp_array_section_kind
+ = c_omp_array_section_kind;
bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
- c_omp_array_section_kind = false;
+ c_omp_array_section_kind = OMP_ARRAY_SECTION_NONE;
c_omp_array_shaping_op_p = false;
ret = make_tree_vector ();
@@ -17034,7 +17065,8 @@ c_parser_omp_variable_list (c_parser *parser,
|| kind == OMP_CLAUSE_FROM))
{
location_t loc = c_parser_peek_token (parser)->location;
- unsigned char save_c_omp_array_section_kind = c_omp_array_section_kind;
+ enum c_omp_array_section_kind save_c_omp_array_section_kind
+ = c_omp_array_section_kind;
bool save_c_omp_array_shaping_op_p = c_omp_array_shaping_op_p;
c_omp_array_section_kind = kind == OMP_CLAUSE_MAP
? OMP_ARRAY_SECTION_UNSTRIDED
diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h
index 4a3571df1e9..65b02a094de 100644
--- a/gcc/c/c-tree.h
+++ b/gcc/c/c-tree.h
@@ -795,10 +795,12 @@ extern int in_countof;
extern int in_typeof;
extern int in_generic;
extern bool c_in_omp_for;
-#define OMP_ARRAY_SECTION_NONE 0
-#define OMP_ARRAY_SECTION_UNSTRIDED 1
-#define OMP_ARRAY_SECTION_STRIDED 2
-extern unsigned char c_omp_array_section_kind;
+enum c_omp_array_section_kind : unsigned char {
+ OMP_ARRAY_SECTION_NONE = 0,
+ OMP_ARRAY_SECTION_UNSTRIDED = 1,
+ OMP_ARRAY_SECTION_STRIDED = 2
+};
+extern enum c_omp_array_section_kind c_omp_array_section_kind;
extern bool c_omp_array_shaping_op_p;
extern bool c_omp_has_array_shape_p;
@@ -880,7 +882,7 @@ extern tree build_component_ref (location_t, tree, tree, location_t,
extern tree handle_counted_by_for_component_ref (location_t, tree);
extern tree build_array_ref (location_t, tree, tree);
extern tree build_omp_array_section (location_t, tree, tree, tree, tree);
-extern tree create_omp_arrayshape_type (tree expr,
+extern tree create_omp_arrayshape_type (location_t loc, tree expr,
vec<tree> *omp_shape_dims);
extern tree build_external_ref (location_t, tree, bool, tree *);
extern void pop_maybe_used (bool);
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 2c696042617..257524fb180 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -84,8 +84,8 @@ int in_generic;
/* True when parsing OpenMP loop expressions. */
bool c_in_omp_for;
-/* True when parsing OpenMP map clause. */
-unsigned char c_omp_array_section_kind;
+/* The kind of OpenMP array section being parsed, if any. */
+enum c_omp_array_section_kind c_omp_array_section_kind;
/* True when parsing OpenMP to/from clause. */
bool c_omp_array_shaping_op_p;
@@ -3872,7 +3872,8 @@ mark_decl_used (tree ref, bool address)
in "myptr[0:2:3][4][5:6]". */
tree
-create_omp_arrayshape_type (tree expr, vec<tree> *omp_shape_dims)
+create_omp_arrayshape_type (location_t loc, tree expr,
+ vec<tree> *omp_shape_dims)
{
tree strip_sections = expr;
@@ -3887,7 +3888,7 @@ create_omp_arrayshape_type (tree expr, vec<tree> *omp_shape_dims)
if (TREE_CODE (type) != POINTER_TYPE)
{
- error ("OpenMP array shaping operator with non-pointer argument");
+ error_at (loc, "OpenMP array shaping operator with non-pointer argument");
return error_mark_node;
}
@@ -3900,7 +3901,7 @@ create_omp_arrayshape_type (tree expr, vec<tree> *omp_shape_dims)
tree maxidx = fold_convert (sizetype, dim);
maxidx = size_binop (MINUS_EXPR, maxidx, size_one_node);
tree index = build_index_type (maxidx);
- type = build_array_type (type, index);
+ type = c_build_array_type (type, index);
}
return type;
@@ -3948,7 +3949,8 @@ build_external_ref (location_t loc, tree id, bool fun, tree *type)
/* For an OpenMP map clause, we can get better diagnostics for decls with
unmappable types if we return the decl with an error_mark_node type,
rather than returning error_mark_node for the decl itself. */
- if (TREE_TYPE (ref) == error_mark_node && !c_omp_array_section_kind)
+ if (TREE_TYPE (ref) == error_mark_node
+ && c_omp_array_section_kind == OMP_ARRAY_SECTION_NONE)
return error_mark_node;
if (TREE_UNAVAILABLE (ref))
@@ -15874,7 +15876,7 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses)
T current expression (initially OMP_CLAUSE_DECL), which is either
a TREE_LIST for array-section-subscript (TREE_PURPOSE is low-bound
expression if specified, TREE_VALUE length expression if specified,
- TREE_CHAIN is what it has been specified after, or some decl.
+ TREE_CHAIN is what it has been specified after), or some decl.
TYPES vector is populated with array section types, MAYBE_ZERO_LEN
set to true if any of the array-section-subscript could have length
of zero (explicit or implicit), FIRST_NON_ONE is the index of the
@@ -15887,7 +15889,14 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses)
<= FIRST_NON_ONE we diagnose non-contiguous arrays if low bound isn't
0 or length isn't the array domain max + 1, for > FIRST_NON_ONE we
can if MAYBE_ZERO_LEN is false. MAYBE_ZERO_LEN will be true in the above
- case though, as some lengths could be zero. */
+ case though, as some lengths could be zero.
+ On entry *DISCONTIGUOUS is 0 if the array section must be contiguous, and
+ non-zero if a discontiguous section is permitted; this function sets
+ *DISCONTIGUOUS to 2 when it determines the section is definitely
+ discontiguous, and otherwise leaves it unchanged.
+ Return the array reference or array-section tree built for T, or T itself
+ (possibly converted to an rvalue) if it is not an array section, or
+ ERROR_MARK_NODE on error. */
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
@@ -16245,11 +16254,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
d = TREE_OPERAND (d, 0))
{
tree d_length = TREE_OPERAND (d, 2);
- tree d_stride = TREE_OPERAND (d, 3);
- if (d_length == NULL_TREE || !integer_onep (d_length)
- || (d_stride && !integer_onep (d_stride)))
+ if (d_length == NULL_TREE || !integer_onep (d_length))
{
- if (discontiguous && *discontiguous)
+ if (*discontiguous)
*discontiguous = 2;
else
{
@@ -16277,11 +16284,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
TREE_OPERAND (t, 1) = lb;
low_bound = lb;
}
- /* NOTE: Stride/length are discarded for affinity/depend here. */
- if (discontiguous
- && *discontiguous
- && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
- && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND)
+ if (*discontiguous)
ret = build_omp_array_section (OMP_CLAUSE_LOCATION (c), ret, low_bound,
length, stride);
else
@@ -16305,26 +16308,31 @@ omp_array_section_low_bound (location_t loc, tree node)
return node;
}
-/* Handle array sections for clause C. On entry *DISCONTIGUOUS is 0 if array
- section must be contiguous, 1 if it can be discontiguous, and in the latter
- case it is set to 2 on exit if it is determined to be discontiguous during
- the function's execution. */
+/* Handle array sections for clause C. DISCONTIGUOUS_P is NULL, or points to
+ a bool that is false, if the array section must be contiguous. If
+ *DISCONTIGUOUS_P is true on entry, a discontiguous section is permitted;
+ on exit it is left true only if the section was determined to be
+ (possibly) discontiguous, and reset to false if it turned out to be
+ contiguous after all. Return true on error. */
static bool
handle_omp_array_sections (tree &c, enum c_omp_region_type ort,
- int *discontiguous = NULL)
+ bool *discontiguous_p = NULL)
{
bool maybe_zero_len = false;
unsigned int first_non_one = 0;
auto_vec<tree, 10> types;
tree *tp = &OMP_CLAUSE_DECL (c);
+ int discontiguous = discontiguous_p && *discontiguous_p ? 1 : 0;
+ if (discontiguous_p)
+ *discontiguous_p = false;
if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY)
&& OMP_ITERATOR_DECL_P (*tp))
tp = &TREE_VALUE (*tp);
tree first = handle_omp_array_sections_1 (c, *tp, types,
maybe_zero_len, first_non_one,
- ort, discontiguous);
+ ort, &discontiguous);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -16392,7 +16400,7 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort,
if (stride == NULL_TREE)
stride = size_one_node;
- if (discontiguous && *discontiguous)
+ if (discontiguous)
{
/* This condition is similar to the error check below, but
whereas that checks for a definitely-discontiguous array
@@ -16418,13 +16426,19 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort,
full_span = true;
}
- if (!integer_onep (stride)
+ /* A stride other than one only matters when the length is
+ also other than one; a length-one dimension touches a
+ single element regardless of stride. */
+ bool length_is_one = length != NULL_TREE && integer_onep (length);
+ bool strided = !integer_onep (stride) && !length_is_one;
+
+ if (strided
|| (higher_discontiguous
&& (!integer_zerop (low_bound)
|| !full_span)))
- *discontiguous = 2;
+ discontiguous = 2;
- if (!integer_onep (stride)
+ if (strided
|| !integer_zerop (low_bound)
|| !full_span)
higher_discontiguous = true;
@@ -16447,8 +16461,8 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort,
if (!tree_int_cst_equal (length, size))
{
is_noncontiguous:
- if (discontiguous && *discontiguous)
- *discontiguous = 2;
+ if (discontiguous)
+ discontiguous = 2;
else
{
error_at (OMP_CLAUSE_LOCATION (c),
@@ -16568,7 +16582,9 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort,
OMP_CLAUSE_DECL (c) = t;
return false;
}
- if (discontiguous && *discontiguous != 2)
+ if (discontiguous_p)
+ *discontiguous_p = discontiguous == 2;
+ if (discontiguous == 1)
first = omp_array_section_low_bound (OMP_CLAUSE_LOCATION (c), first);
first = c_fully_fold (first, false, NULL);
OMP_CLAUSE_DECL (c) = first;
@@ -16585,8 +16601,9 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort,
OMP_CLAUSE_SIZE (c) = size;
}
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
- && !(discontiguous && *discontiguous == 2))
+ /* Done, except for MAP clauses and for array sections that may be
+ noncontiguous. */
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP && discontiguous != 2)
return false;
auto_vec<omp_addr_token *, 10> addr_tokens;
@@ -16617,7 +16634,7 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort,
&& addr_tokens[0]->u.structure_base_kind == BASE_DECL
&& addr_tokens[1]->type == ACCESS_METHOD
&& omp_access_chain_p (addr_tokens, 1))
- || (discontiguous && *discontiguous == 2))
+ || discontiguous == 2)
c = nc;
return false;
@@ -17868,7 +17885,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
grp_start_p = pc;
grp_sentinel = OMP_CLAUSE_CHAIN (c);
- int discontiguous
+ bool discontiguous
= (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM);
if (handle_omp_array_sections (c, ort, &discontiguous))
@@ -17884,7 +17901,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
pre-existing "chained access" case, where C is
reassigned only to advance past appended nodes that
stay chained after the original, unmoved clause.) */
- if (discontiguous == 2)
+ if (discontiguous)
*pc = c;
t = OMP_CLAUSE_DECL (c);
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-1.c b/gcc/testsuite/c-c++-common/gomp/array-section-1.c
new file mode 100644
index 00000000000..c0561284ac4
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-1.c
@@ -0,0 +1,20 @@
+/* { dg-do compile } */
+
+/* Check that array section syntax with empty length is handled correctly.
+ '::1' is valid only in C. */
+
+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: : ])
+
+#pragma omp target update to(A[:::n], B[:::n][: :: n : ])
+/* { dg-error "expected expression before ':' token" "" { target c } .-1 } */
+/* { dg-error "expected expression before '::' token" "" { target c } .-2 } */
+/* { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target c } .-3 } */
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-2.c b/gcc/testsuite/c-c++-common/gomp/array-section-2.c
new file mode 100644
index 00000000000..4daa0a6b497
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-2.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 c++ } .-2 }
+ // { dg-error "expected an OpenMP clause before '\\\]' token" "" { target c++ } .-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 c++ } .-2 }
+ // { dg-error "expected an OpenMP clause before '\\\]' token" "" { target c++ } .-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/c-c++-common/gomp/bad-array-shaping-1.c b/gcc/testsuite/c-c++-common/gomp/bad-array-shaping-1.c
new file mode 100644
index 00000000000..fafbfc86086
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/bad-array-shaping-1.c
@@ -0,0 +1,12 @@
+/* { dg-do compile } */
+
+/* The OpenMP array-shaping operator requires a pointer argument. */
+
+int x;
+
+void f (void)
+{
+#pragma omp target update to( ([10])x )
+/* { dg-error "OpenMP array shaping operator with non-pointer argument" "" { target *-*-* } .-1 } */
+/* { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target c++ } .-2 } */
+}
diff --git a/gcc/testsuite/gcc.dg/gomp/array-shaping-1.c b/gcc/testsuite/gcc.dg/gomp/array-shaping-1.c
new file mode 100644
index 00000000000..9bb08acf0ae
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/array-shaping-1.c
@@ -0,0 +1,17 @@
+// { dg-do compile }
+
+/* The C front end used to ICE building the type for a variable-bound
+ dimension of the OpenMP array-shaping operator whenever the
+ resulting shaped expression was then further subscripted with an array
+ section. create_omp_arrayshape_type built each dimension with a
+ build_array_type call instead of *c_*build_array_type, leaving the
+ C-specific "variable size" bits inconsistent for later reuse. */
+
+void
+array_shaping (double *a, int nx, int ny)
+{
+#pragma omp target update from( (([nx][ny + 2])a)[0:nx][1], \
+ (([nx][ny + 2])a)[0:nx][ny] )
+#pragma omp target update to( (([nx][ny + 2])a)[0:nx][0], \
+ (([nx][ny + 2])a)[0:nx][ny + 1] )
+}
--
2.55.0
