[PATCH 2/5] openmp: Support pointer indirection in array-shaping casts and array sections
Paul-Antoine Arras <[email protected]> Wed, 5 Aug 2026 22:23:40 +0200
| Newsgroups | gmane.comp.gcc.patches |
|---|---|
| Message-ID | <[email protected]> |
Array-shaping casts and noncontiguous array sections were already
partially supported; what is new here is handling further pointer
indirection within them, i.e. array-of-pointers sections and
array-shaping casts that cross more than one pointer. Each further
indirection starts a new "segment" of dimensions, recorded by marking
the GOMP_MAP_GRID_DIM clause for that dimension with
OMP_CLAUSE_MAP_GRID_DIM_POINTER; a GOMP_MAP_SHAPE_DIM clause is emitted
alongside dimensions introduced by an array-shaping cast so the
lowering pass can size them.
This also fixes crashes/leaks uncovered along the way (member access
via a dependent object, 2+-dimensional array-of-pointers in templates),
rejects array-shaping casts over a range of pointers per the OpenMP
spec, accepts decayed array parameters, and teaches
omp_parse_noncontiguous_array to look through the VIEW_CONVERT_EXPR an
array-shaping cast introduces over an OMP_ARRAY_SECTION.
Add tests scanning the "original" dump to check the clause chains built
for each case; the corresponding "lower" dump checks for the generated
descriptors are added once the omp-low.cc lowering support lands.
gcc/c-family/ChangeLog:
* c-omp.cc (omp_expand_grid_dim): Support array-of-pointers sections
spanning multiple segments and dimensions introduced by an
array-shaping cast, marking dimensions that select through a pointer
and emitting GOMP_MAP_SHAPE_DIM clauses for shape-cast dimensions.
(omp_handle_noncontig_array): Adjust to omp_expand_grid_dim's new
signature and synthesize whole-span GOMP_MAP_GRID_DIM/
GOMP_MAP_SHAPE_DIM clauses for any shape-cast dimensions left over.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_variable_list): Accept array-shaping
casts over decayed array parameters.
* c-typeck.cc (create_omp_arrayshape_type): Reject array-shaping
casts over a range of pointers.
(handle_omp_array_sections_1): Support array-of-pointers sections.
(handle_omp_array_sections): Likewise.
(c_finish_omp_clauses): Likewise.
gcc/cp/ChangeLog:
* decl.cc (cp_omp_create_arrayshape_type): Reject array-shaping casts
over a range of pointers.
* parser.cc (cp_parser_omp_var_list_no_open): Accept array-shaping
casts over decayed array parameters.
* semantics.cc (handle_omp_array_sections_1): Support
array-of-pointers sections; fix crash on array-shaping cast of a
member via a dependent object and leaks for 2+-dimensional
array-of-pointers in templates.
(finish_omp_clauses): Support array-of-pointers sections.
gcc/ChangeLog:
* omp-general.cc (omp_parse_noncontiguous_array): Look through the
VIEW_CONVERT_EXPR an array-shaping cast introduces over an
OMP_ARRAY_SECTION.
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/array-section-3.c: New test.
* c-c++-common/gomp/array-section-4.c: New test.
* c-c++-common/gomp/array-section-5.c: New test.
* c-c++-common/gomp/array-section-6.c: New test.
* c-c++-common/gomp/array-section-8.c: New test.
* c-c++-common/gomp/array-section-9.c: New test.
* c-c++-common/gomp/target-update-iterators-4.c: New test.
* g++.dg/gomp/array-shaping-3.C: New test.
* g++.dg/gomp/array-shaping-4.C: New test.
* g++.dg/gomp/bad-array-shaping-9.C: New test.
* gcc.dg/gomp/array-shaping-9.c: New test.
* gcc.dg/gomp/bad-array-section-c-9.c: New test.
* gcc.dg/gomp/bad-array-shaping-c-8.c: New test.
---
gcc/c-family/c-omp.cc | 151 ++++++++++++++----
gcc/c/c-parser.cc | 59 ++++++-
gcc/c/c-typeck.cc | 144 ++++++++++++-----
gcc/cp/decl.cc | 50 +++++-
gcc/cp/parser.cc | 69 +++++++-
gcc/cp/semantics.cc | 114 ++++++++-----
gcc/omp-general.cc | 5 +-
.../c-c++-common/gomp/array-section-1.c | 33 ++++
.../c-c++-common/gomp/array-section-2.c | 15 ++
.../c-c++-common/gomp/array-section-3.c | 38 +++++
.../c-c++-common/gomp/array-section-4.c | 21 +++
.../c-c++-common/gomp/array-section-5.c | 25 +++
.../c-c++-common/gomp/array-section-6.c | 18 +++
.../c-c++-common/gomp/array-section-8.c | 27 ++++
.../c-c++-common/gomp/array-section-9.c | 21 +++
.../gomp/target-update-iterators-4.c | 19 +++
gcc/testsuite/g++.dg/gomp/array-shaping-3.C | 71 ++++++++
gcc/testsuite/g++.dg/gomp/array-shaping-4.C | 37 +++++
.../g++.dg/gomp/bad-array-shaping-9.C | 51 ++++++
gcc/testsuite/gcc.dg/gomp/array-shaping-9.c | 25 +++
.../gcc.dg/gomp/bad-array-section-c-9.c | 60 +++++++
.../gcc.dg/gomp/bad-array-shaping-c-8.c | 45 ++++++
22 files changed, 973 insertions(+), 125 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/array-section-3.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-4.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-5.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-6.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-8.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-9.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c
create mode 100644 gcc/testsuite/g++.dg/gomp/array-shaping-3.C
create mode 100644 gcc/testsuite/g++.dg/gomp/array-shaping-4.C
create mode 100644 gcc/testsuite/g++.dg/gomp/bad-array-shaping-9.C
create mode 100644 gcc/testsuite/gcc.dg/gomp/array-shaping-9.c
create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-section-c-9.c
create mode 100644 gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-8.c
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 01721cafad0..c4941166206 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -3718,60 +3718,110 @@ omp_expand_access_chain (tree *pc, tree expr,
return pc;
}
-static tree *
-omp_expand_grid_dim (location_t loc, tree *pc, tree decl)
-{
- if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
- pc = omp_expand_grid_dim (loc, pc, TREE_OPERAND (decl, 0));
- else
- return pc;
+/* Append a GOMP_MAP_GRID_DIM (and possibly GOMP_MAP_GRID_STRIDE and
+ GOMP_MAP_SHAPE_DIM) clause for each dimension of the OMP_ARRAY_SECTION chain
+ in DECL onto the clause chain at *PC, base to outer. PC points to the
+ insertion point in the clause chain. DECL is the OMP_ARRAY_SECTION being
+ expanded, or the base declaration once the chain is exhausted. *TYPE is set
+ to the base declaration's type, then peeled through recursion so a dimension
+ is marked OMP_CLAUSE_MAP_GRID_DIM_POINTER when it is a pointer rather than an
+ array (output only). *FIRST is true only for the first dimension, so that it
+ is never marked as pointer even when the base type has decayed (output only).
+ *SHAPE_TYPE is set to the array-shaping type for the GOMP_MAP_SHAPE_DIM
+ clause (output only). Returns the updated chain insertion point. */
- tree c = *pc;
+static tree *
+omp_expand_grid_dim (location_t loc, tree *pc, tree decl, tree *type,
+ bool *first, tree *shape_type)
+{
+ if (TREE_CODE (decl) == OMP_ARRAY_SECTION || TREE_CODE (decl) == ARRAY_REF)
+ {
+ tree op = TREE_OPERAND (decl, 0);
+ if (TREE_CODE (op) == VIEW_CONVERT_EXPR
+ && TREE_CODE (TREE_OPERAND (op, 0)) == OMP_ARRAY_SECTION)
+ {
+ /* Array section within an array-shaping cast. */
+ pc = omp_expand_grid_dim (loc, pc, TREE_OPERAND (op, 0), type,
+ first, shape_type);
+ *shape_type = TREE_TYPE (op);
+ }
+ else
+ pc = omp_expand_grid_dim (loc, pc, op, type, first, shape_type);
+ }
+ else
+ {
+ *type = TREE_TYPE (decl);
+ *first = true;
+ *shape_type = NULL_TREE;
+ return pc;
+ }
+
+ tree nc = OMP_CLAUSE_CHAIN (*pc);
tree low_bound = TREE_OPERAND (decl, 1);
tree length = TREE_OPERAND (decl, 2);
tree stride = TREE_OPERAND (decl, 3);
- tree cd = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (cd, GOMP_MAP_GRID_DIM);
- OMP_CLAUSE_DECL (cd) = unshare_expr (low_bound);
- OMP_CLAUSE_SIZE (cd) = unshare_expr (length);
+ tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_GRID_DIM);
+ OMP_CLAUSE_DECL (c) = unshare_expr (low_bound);
+ OMP_CLAUSE_SIZE (c) = length ? unshare_expr (length) : size_one_node;
+ if (!*first && TREE_CODE (*type) == POINTER_TYPE)
+ OMP_CLAUSE_MAP_GRID_DIM_POINTER (c) = 1;
+ OMP_CLAUSE_CHAIN (*pc) = c;
+ pc = &OMP_CLAUSE_CHAIN (*pc);
if (stride && !integer_onep (stride))
{
- tree cs = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (cs, GOMP_MAP_GRID_STRIDE);
- OMP_CLAUSE_DECL (cs) = unshare_expr (stride);
-
- OMP_CLAUSE_CHAIN (cs) = OMP_CLAUSE_CHAIN (c);
- OMP_CLAUSE_CHAIN (cd) = cs;
- OMP_CLAUSE_CHAIN (c) = cd;
- pc = &OMP_CLAUSE_CHAIN (cd);
+ c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_GRID_STRIDE);
+ OMP_CLAUSE_DECL (c) = unshare_expr (stride);
+ OMP_CLAUSE_CHAIN (*pc) = c;
+ pc = &OMP_CLAUSE_CHAIN (*pc);
}
- else
+ if (*shape_type)
{
- OMP_CLAUSE_CHAIN (cd) = OMP_CLAUSE_CHAIN (c);
- OMP_CLAUSE_CHAIN (c) = cd;
- pc = &OMP_CLAUSE_CHAIN (c);
- }
+ c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_SHAPE_DIM);
+ tree dtype = TYPE_DOMAIN (*shape_type);
+ tree minval = TYPE_MIN_VALUE (dtype);
+ tree maxval = TYPE_MAX_VALUE (dtype);
+ minval = fold_convert (sizetype, minval);
+ maxval = fold_convert (sizetype, maxval);
+ tree dim = size_binop (MINUS_EXPR, maxval, minval);
+ dim = size_binop (PLUS_EXPR, dim, size_one_node);
+ OMP_CLAUSE_DECL (c) = dim;
+ OMP_CLAUSE_CHAIN (*pc) = c;
+ pc = &OMP_CLAUSE_CHAIN (*pc);
+ tree elt = TREE_TYPE (*shape_type);
+ *shape_type = TREE_CODE (elt) == ARRAY_TYPE ? elt : NULL_TREE;
+ }
+ OMP_CLAUSE_CHAIN (c) = nc;
+
+ *first = false;
+ *type = *shape_type ? *shape_type : TREE_TYPE (*type);
return pc;
}
+/* Replace clause C, a TO/FROM clause on a non-contiguous array section, at
+ *PC with a GOMP_MAP_TO_GRID/GOMP_MAP_FROM_GRID clause describing BASE
+ (the section's base declaration), followed by the GOMP_MAP_GRID_DIM/
+ GOMP_MAP_GRID_STRIDE clauses for each dimension of C's OMP_ARRAY_SECTION
+ chain. Returns the updated chain insertion point. */
+
tree *
omp_handle_noncontig_array (location_t loc, tree *pc, tree c, tree base)
{
- tree type;
+ tree eltype = TREE_TYPE (base);
- if (POINTER_TYPE_P (TREE_TYPE (base)))
- type = TREE_TYPE (TREE_TYPE (base));
- else
- type = strip_array_types (TREE_TYPE (base));
+ while (TREE_CODE (eltype) == ARRAY_TYPE || POINTER_TYPE_P (eltype))
+ eltype = TREE_TYPE (eltype);
tree c_map = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_DECL (c_map) = unshare_expr (base);
/* Use the element size (or pointed-to type size) here. */
- OMP_CLAUSE_SIZE (c_map) = TYPE_SIZE_UNIT (type);
+ OMP_CLAUSE_SIZE (c_map) = TYPE_SIZE_UNIT (eltype);
switch (OMP_CLAUSE_CODE (c))
{
@@ -3789,7 +3839,44 @@ omp_handle_noncontig_array (location_t loc, tree *pc, tree c, tree base)
*pc = c_map;
- return omp_expand_grid_dim (loc, pc, OMP_CLAUSE_DECL (c));
+ tree dtype;
+ bool first;
+ tree shape_type;
+ pc = omp_expand_grid_dim (loc, pc, OMP_CLAUSE_DECL (c), &dtype, &first,
+ &shape_type);
+
+ /* Build whole-span dimensions for remaining shape_type layers. */
+ tree nc = OMP_CLAUSE_CHAIN (*pc);
+ tree dc = NULL_TREE;
+ while (shape_type != NULL_TREE && TREE_CODE (shape_type) == ARRAY_TYPE)
+ {
+ /* Compute whole-span length. */
+ tree dtype = TYPE_DOMAIN (shape_type);
+ gcc_assert (integer_zerop (TYPE_MIN_VALUE (dtype)));
+ tree length = TYPE_MAX_VALUE (dtype);
+ length = size_binop (PLUS_EXPR, length, size_one_node);
+
+ /* Build grid_dim node. */
+ dc = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (dc, GOMP_MAP_GRID_DIM);
+ OMP_CLAUSE_DECL (dc) = size_zero_node;
+ OMP_CLAUSE_SIZE (dc) = length;
+ OMP_CLAUSE_CHAIN (*pc) = dc;
+ pc = &OMP_CLAUSE_CHAIN (*pc);
+
+ /* Build shape_dim node. */
+ dc = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (dc, GOMP_MAP_SHAPE_DIM);
+ OMP_CLAUSE_DECL (dc) = length;
+ OMP_CLAUSE_CHAIN (*pc) = dc;
+ pc = &OMP_CLAUSE_CHAIN (*pc);
+
+ shape_type = TREE_TYPE (shape_type);
+ }
+ if (dc != NULL_TREE)
+ OMP_CLAUSE_CHAIN (dc) = nc;
+
+ return pc;
}
/* Translate "array_base_decl access_method" to OMP mapping clauses. */
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 1c54af9910e..3d5a4825a10 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -17017,6 +17017,18 @@ c_parser_omp_variable_list (c_parser *parser,
sections--;
}
+ /* The previous loop may have uncovered a shaping cast. DECL_P
+ excludes such a cast over a further array section (reshaping
+ e.g. "x[1][1:2]"): that case is left for omp_expand_grid_dim's
+ SHAPE_TYPE handling in c-omp.cc. */
+ if (!reshaped_to && TREE_CODE (decl) == VIEW_CONVERT_EXPR
+ && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
+ && DECL_P (TREE_OPERAND (decl, 0)))
+ {
+ reshaped_to = TREE_TYPE (decl);
+ decl = TREE_OPERAND (decl, 0);
+ }
+
/* The handling of INDIRECT_REF here in the presence of
array-shaping operations is a little tricky. We need to
avoid treating a pointer dereference as a unit-sized array
@@ -17087,7 +17099,52 @@ c_parser_omp_variable_list (c_parser *parser,
}
}
- if (reshaped_to)
+ if (reshaped_to && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+ {
+ unsigned array_layers = 0;
+ tree basetype = TREE_TYPE (decl);
+ while (TREE_CODE (basetype) == ARRAY_TYPE)
+ {
+ basetype = TREE_TYPE (basetype);
+ array_layers++;
+ }
+
+ /* A further level of indirection here (e.g. "int **x[N]")
+ is not supported (for now); treating the second pointer's
+ bytes as the reshaped array's elements would silently read
+ the wrong data. (See also
+ gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-8.c). */
+ if (TREE_CODE (TREE_TYPE (basetype)) == POINTER_TYPE)
+ {
+ sorry_at (loc, "more than one level of pointer "
+ "indirection in a noncontiguous target "
+ "update");
+ decl = error_mark_node;
+ }
+ else if (dims.length () != array_layers)
+ {
+ error_at (loc, "too many array section specifiers "
+ "for %qT", reshaped_to);
+ decl = error_mark_node;
+ }
+ else
+ {
+ for (unsigned i = 0; i < array_layers; i++)
+ {
+ omp_dim &d = dims[dims.length () - 1 - i];
+ decl = build_omp_array_section (loc, decl,
+ d.low_bound, d.length,
+ d.stride);
+ }
+ dims.truncate (0);
+ decl = build1_loc (loc, VIEW_CONVERT_EXPR, reshaped_to,
+ decl);
+ tree elems = c_array_type_nelts_total (reshaped_to);
+ decl = build_omp_array_section (loc, decl, size_zero_node,
+ elems, NULL_TREE);
+ }
+ }
+ else if (reshaped_to)
{
unsigned reshaped_dims = 0;
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 0c141520501..ff061044ee4 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -3877,13 +3877,35 @@ create_omp_arrayshape_type (tree expr, vec<tree> *omp_shape_dims)
while (TREE_CODE (strip_sections) == OMP_ARRAY_SECTION
|| TREE_CODE (strip_sections) == ARRAY_REF)
- strip_sections = TREE_OPERAND (strip_sections, 0);
+ {
+ /* The array-shaping operator only applies to a single pointer, not an
+ array section. */
+ if (TREE_CODE (strip_sections) == OMP_ARRAY_SECTION)
+ {
+ tree array = TREE_OPERAND (strip_sections, 0);
+ tree length = TREE_OPERAND (strip_sections, 2);
+ tree atype = TREE_TYPE (array);
+ tree eltype = atype != NULL_TREE ? TREE_TYPE (atype) : NULL_TREE;
+ if (eltype != NULL_TREE
+ && POINTER_TYPE_P (eltype)
+ && (length == NULL_TREE || !integer_onep (length)))
+ {
+ error ("OpenMP array shaping operator with non-pointer "
+ "argument");
+ return error_mark_node;
+ }
+ }
+ strip_sections = TREE_OPERAND (strip_sections, 0);
+ }
tree type = TREE_TYPE (strip_sections);
if (TREE_CODE (type) == REFERENCE_TYPE)
type = TREE_TYPE (type);
+ while (TREE_CODE (type) == ARRAY_TYPE)
+ type = TREE_TYPE (type);
+
if (TREE_CODE (type) != POINTER_TYPE)
{
error ("OpenMP array shaping operator with non-pointer argument");
@@ -15864,14 +15886,16 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses)
can if MAYBE_ZERO_LEN is false. MAYBE_ZERO_LEN will be true in the above
case though, as some lengths could be zero.
NON_CONTIGUOUS will be true if this is an OpenACC non-contiguous array
- section. */
+ section. TYPE carries the type of the inner section through recursion
+ (output only). */
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
bool &maybe_zero_len, unsigned int &first_non_one,
- bool &non_contiguous, enum c_omp_region_type ort, int *discontiguous)
+ bool &non_contiguous, enum c_omp_region_type ort,
+ int *discontiguous, tree *type)
{
- tree ret, low_bound, length, stride, type;
+ tree ret, low_bound, length, stride;
bool openacc = (ort & C_ORT_ACC) != 0;
if (TREE_CODE (t) != OMP_ARRAY_SECTION)
{
@@ -15941,20 +15965,16 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
ret = convert_lvalue_to_rvalue (OMP_CLAUSE_LOCATION (c),
ret, false, false);
}
+ *type = TREE_TYPE (ret);
return ret;
}
ret = handle_omp_array_sections_1 (c, TREE_OPERAND (t, 0), types,
maybe_zero_len, first_non_one,
- non_contiguous, ort,
- discontiguous);
+ non_contiguous, ort, discontiguous, type);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
- if (TREE_CODE (ret) == OMP_ARRAY_SECTION)
- type = TREE_TYPE (TREE_TYPE (TREE_OPERAND (ret, 0)));
- else
- type = TREE_TYPE (ret);
low_bound = TREE_OPERAND (t, 1);
length = TREE_OPERAND (t, 2);
stride = TREE_OPERAND (t, 3);
@@ -16041,11 +16061,11 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
&& (TREE_CODE (length) != INTEGER_CST || integer_onep (length)))
first_non_one++;
}
- if (TREE_CODE (type) == ARRAY_TYPE)
+ if (TREE_CODE (*type) == ARRAY_TYPE)
{
if (length == NULL_TREE
- && (TYPE_DOMAIN (type) == NULL_TREE
- || TYPE_MAX_VALUE (TYPE_DOMAIN (type)) == NULL_TREE))
+ && (TYPE_DOMAIN (*type) == NULL_TREE
+ || TYPE_MAX_VALUE (TYPE_DOMAIN (*type)) == NULL_TREE))
{
error_at (OMP_CLAUSE_LOCATION (c),
"for unknown bound array type length expression must "
@@ -16069,13 +16089,11 @@ 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 (TYPE_DOMAIN (type)
- && TYPE_MAX_VALUE (TYPE_DOMAIN (type))
- && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (type)))
- == INTEGER_CST)
+ if (TYPE_DOMAIN (*type) && TYPE_MAX_VALUE (TYPE_DOMAIN (*type))
+ && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (*type))) == INTEGER_CST)
{
tree size
- = fold_convert (sizetype, TYPE_MAX_VALUE (TYPE_DOMAIN (type)));
+ = fold_convert (sizetype, TYPE_MAX_VALUE (TYPE_DOMAIN (*type)));
size = size_binop (PLUS_EXPR, size, size_one_node);
if (TREE_CODE (low_bound) == INTEGER_CST)
{
@@ -16102,11 +16120,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
}
maybe_zero_len = true;
}
- else if (length == NULL_TREE
- && first_non_one == types.length ()
- && tree_int_cst_equal
- (TYPE_MAX_VALUE (TYPE_DOMAIN (type)),
- low_bound))
+ else if (length == NULL_TREE && first_non_one == types.length ()
+ && tree_int_cst_equal (
+ TYPE_MAX_VALUE (TYPE_DOMAIN (*type)), low_bound))
first_non_one++;
}
else if (length == NULL_TREE)
@@ -16188,7 +16204,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
}
}
}
- else if (TREE_CODE (type) == POINTER_TYPE)
+ else if (TREE_CODE (*type) == POINTER_TYPE)
{
if (length == NULL_TREE)
{
@@ -16267,7 +16283,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND)
- types.safe_push (type);
+ types.safe_push (*type);
/* We will need to evaluate lb more than once. */
tree lb = save_expr (low_bound);
if (lb != low_bound)
@@ -16284,6 +16300,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
length, stride);
else
ret = build_array_ref (OMP_CLAUSE_LOCATION (c), ret, low_bound);
+
+ *type = TREE_TYPE (*type);
return ret;
}
@@ -16322,9 +16340,10 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort,
|| 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,
- non_contiguous, ort, discontiguous);
+ tree type;
+ tree first
+ = handle_omp_array_sections_1 (c, *tp, types, maybe_zero_len, first_non_one,
+ non_contiguous, ort, discontiguous, &type);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -16343,8 +16362,8 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort,
if (tem == NULL_TREE)
tem = TREE_VALUE (t);
else
- tem = build2 (COMPOUND_EXPR, TREE_TYPE (tem),
- TREE_VALUE (t), tem);
+ tem = build2 (COMPOUND_EXPR, TREE_TYPE (tem), TREE_VALUE (t),
+ tem);
}
t = TREE_CHAIN (t);
}
@@ -16364,6 +16383,11 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort,
maybe_zero_len = true;
bool higher_discontiguous = false;
+ /* Whether the first dimension uses the extended array section syntax,
+ i.e. whether it specifies a length or a stride. */
+ bool first_extended = false;
+ /* Whether any dimension container except the last is a pointer. */
+ bool non_last_pointer = false;
for (i = num, t = OMP_CLAUSE_DECL (c); i > 0;
t = TREE_OPERAND (t, 0))
@@ -16427,10 +16451,15 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort,
full_span = true;
}
+ tree op = TREE_OPERAND (t, 0);
+ bool shaped
+ = TREE_CODE (op) == VIEW_CONVERT_EXPR
+ && TREE_CODE (TREE_TYPE (op)) == ARRAY_TYPE
+ && TREE_CODE (TREE_OPERAND (op, 0)) == OMP_ARRAY_SECTION;
if (!integer_onep (stride)
|| (higher_discontiguous
- && (!integer_zerop (low_bound)
- || !full_span)))
+ && (!integer_zerop (low_bound) || !full_span))
+ || shaped)
*discontiguous = 2;
if (!integer_onep (stride)
@@ -16440,20 +16469,29 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort,
}
if (!maybe_zero_len && i > first_non_one)
{
+ bool full_span;
if (integer_nonzerop (low_bound))
goto is_noncontiguous;
- if (length != NULL_TREE
- && TREE_CODE (length) == INTEGER_CST
- && TYPE_DOMAIN (types[i])
- && TYPE_MAX_VALUE (TYPE_DOMAIN (types[i]))
- && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])))
- == INTEGER_CST)
+ if (length != NULL_TREE && TREE_CODE (length) == INTEGER_CST)
{
- tree size;
- size = size_binop (PLUS_EXPR,
- TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])),
- size_one_node);
- if (!tree_int_cst_equal (length, size))
+ /* A pointer-typed dimension has no static bound to compare
+ against, so it can never be proven to span the whole
+ dimension; treat it conservatively as not full span. */
+ full_span = false;
+ if (TREE_CODE (types[i]) == ARRAY_TYPE
+ && TYPE_DOMAIN (types[i])
+ && TYPE_MAX_VALUE (TYPE_DOMAIN (types[i]))
+ && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])))
+ == INTEGER_CST)
+ {
+ tree size;
+ size
+ = size_binop (PLUS_EXPR,
+ TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])),
+ size_one_node);
+ full_span = tree_int_cst_equal (length, size);
+ }
+ if (!full_span)
{
is_noncontiguous:
if (discontiguous && *discontiguous)
@@ -16538,7 +16576,24 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort,
else
size = size_binop (MULT_EXPR, size, l);
}
+
+ if (i == 0
+ && ((length != NULL_TREE && !integer_onep (length))
+ || !integer_onep (stride)))
+ first_extended = true;
+ if (i < num - 1 && TREE_CODE (types[i]) == POINTER_TYPE
+ && /* accept decayed array */ !(
+ i == 0 && TREE_CODE (TREE_TYPE (types[i])) == ARRAY_TYPE))
+ non_last_pointer = true;
}
+
+ if (first_extended && non_last_pointer)
+ error_at (
+ OMP_CLAUSE_LOCATION (c),
+ "only the last dimension can refer to a pointer when the first "
+ "dimension uses the extended array section syntax in %qs clause",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+
if (non_contiguous)
{
int kind = OMP_CLAUSE_MAP_KIND (c);
@@ -17942,7 +17997,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
}
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE)
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_SHAPE_DIM)
break;
/* FALLTHRU */
case OMP_CLAUSE_TO:
diff --git a/gcc/cp/decl.cc b/gcc/cp/decl.cc
index 1d67999cfef..cb590165091 100644
--- a/gcc/cp/decl.cc
+++ b/gcc/cp/decl.cc
@@ -13645,17 +13645,60 @@ cp_omp_create_arrayshape_type (location_t loc, tree expr,
while (TREE_CODE (strip_sections) == OMP_ARRAY_SECTION
|| TREE_CODE (strip_sections) == ARRAY_REF)
- strip_sections = TREE_OPERAND (strip_sections, 0);
+ {
+ /* The array-shaping operator only applies to a single pointer, not an
+ array section. */
+ if (TREE_CODE (strip_sections) == OMP_ARRAY_SECTION)
+ {
+ tree array = TREE_OPERAND (strip_sections, 0);
+ tree length = TREE_OPERAND (strip_sections, 2);
+ tree atype = TREE_TYPE (array);
+ tree eltype = atype != NULL_TREE ? TREE_TYPE (atype) : NULL_TREE;
+ if (eltype != NULL_TREE
+ && POINTER_TYPE_P (eltype)
+ && (length == NULL_TREE || !integer_onep (length)))
+ {
+ error ("OpenMP array shaping operator with non-pointer "
+ "argument");
+ return error_mark_node;
+ }
+ }
+ strip_sections = TREE_OPERAND (strip_sections, 0);
+ }
/* Determine the element type, either directly or by using
"decltype" of an expression representing an element to
figure it out later during template instantiation. */
if (type_dependent_expression_p (expr))
{
+ /* Peel ARRAY_TYPE layers before deferring the final dereference,
+ mirroring the non-dependent case below. Without this, a single
+ deferred dereference of the un-indexed base only strips one layer
+ via ordinary array-to-pointer decay, leaving any further array
+ dimensions in the resolved type, where they are later mistaken for an
+ omitted trailing shape dimension. (See also
+ gcc/testsuite/g++.dg/gomp/array-shaping-3.C). */
+ tree btype = TREE_TYPE (strip_sections);
+
+ tree base_expr = strip_sections;
+ if (btype != NULL_TREE)
+ {
+ if (TREE_CODE (btype) == REFERENCE_TYPE)
+ btype = TREE_TYPE (btype);
+
+ while (TREE_CODE (btype) == ARRAY_TYPE)
+ {
+ base_expr
+ = build_min_nt_loc (loc, ARRAY_REF, base_expr,
+ integer_zero_node, NULL_TREE, NULL_TREE);
+ btype = TREE_TYPE (btype);
+ }
+ }
+
type = cxx_make_type (DECLTYPE_TYPE);
DECLTYPE_TYPE_EXPR (type)
- = build_min_nt_loc (loc, INDIRECT_REF, strip_sections);
+ = build_min_nt_loc (loc, INDIRECT_REF, base_expr);
DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (type) = true;
SET_TYPE_STRUCTURAL_EQUALITY (type);
}
@@ -13666,6 +13709,9 @@ cp_omp_create_arrayshape_type (location_t loc, tree expr,
if (TREE_CODE (type) == REFERENCE_TYPE)
type = TREE_TYPE (type);
+ while (TREE_CODE (type) == ARRAY_TYPE)
+ type = TREE_TYPE (type);
+
if (TREE_CODE (type) != POINTER_TYPE)
{
error ("OpenMP array shaping operator with non-pointer argument");
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index c0c3759e6da..2097cd3ffd1 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -41635,9 +41635,10 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
location_t loc = token->location;
decl = cp_parser_assignment_expression (parser);
- if ((TREE_CODE (decl) == VIEW_CONVERT_EXPR
- && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
- || TREE_CODE (decl) == OMP_ARRAYSHAPE_CAST_EXPR)
+ if (TREE_TYPE (decl) != NULL_TREE
+ && ((TREE_CODE (decl) == VIEW_CONVERT_EXPR
+ && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+ || TREE_CODE (decl) == OMP_ARRAYSHAPE_CAST_EXPR))
{
reshaped_to = TREE_TYPE (decl);
decl = TREE_OPERAND (decl, 0);
@@ -41690,6 +41691,19 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
sections--;
}
+ /* The previous loop may have uncovered a shaping cast. DECL_P
+ excludes such a cast over a further array section (reshaping
+ e.g. "x[1][1:2]"): that case is left for omp_expand_grid_dim's
+ SHAPE_TYPE handling in c-omp.cc. */
+ if (!reshaped_to && TREE_TYPE (decl) != NULL_TREE
+ && TREE_CODE (decl) == VIEW_CONVERT_EXPR
+ && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
+ && DECL_P (TREE_OPERAND (decl, 0)))
+ {
+ reshaped_to = TREE_TYPE (decl);
+ decl = TREE_OPERAND (decl, 0);
+ }
+
/* The handling of INDIRECT_REF here in the presence of
array-shaping operations is a little tricky. We need to
avoid treating a pointer dereference as a unit-sized array
@@ -41764,7 +41778,54 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
}
}
- if (reshaped_to)
+ if (reshaped_to && TREE_TYPE (decl) != NULL_TREE
+ && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+ {
+ unsigned array_layers = 0;
+ tree basetype = TREE_TYPE (decl);
+ while (TREE_CODE (basetype) == ARRAY_TYPE)
+ {
+ basetype = TREE_TYPE (basetype);
+ array_layers++;
+ }
+
+ /* A further level of indirection here (e.g. "int **x[N]")
+ is not supported (for now); treating the second pointer's
+ bytes as the reshaped array's elements would silently read
+ the wrong data. (See also
+ gcc/testsuite/g++.dg/gomp/bad-array-shaping-9.C). */
+ if (TREE_CODE (TREE_TYPE (basetype)) == POINTER_TYPE)
+ {
+ sorry_at (loc, "more than one level of pointer "
+ "indirection in a noncontiguous target "
+ "update");
+ decl = error_mark_node;
+ }
+ else if (dims.length () != array_layers)
+ {
+ error_at (loc, "too many array section specifiers "
+ "for %qT", reshaped_to);
+ decl = error_mark_node;
+ }
+ else
+ {
+ for (unsigned i = 0; i < array_layers; i++)
+ {
+ omp_dim &d = dims[dims.length () - 1 - i];
+ decl = grok_omp_array_section (loc, decl,
+ d.low_bound, d.length,
+ d.stride);
+ }
+ dims.truncate (0);
+ decl = cp_build_omp_arrayshape_cast (loc, reshaped_to,
+ decl,
+ tf_warning_or_error);
+ tree elems = array_type_nelts_total (reshaped_to);
+ decl = grok_omp_array_section (loc, decl, size_zero_node,
+ elems, NULL_TREE);
+ }
+ }
+ else if (reshaped_to)
{
unsigned reshaped_dims = 0;
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 4bf65d28fc9..16abf6b4b56 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -6053,15 +6053,16 @@ public:
can if MAYBE_ZERO_LEN is false. MAYBE_ZERO_LEN will be true in the above
case though, as some lengths could be zero.
NON_CONTIGUOUS will be true if this is an OpenACC non-contiguous array
- section. */
+ section. TYPE carries the type of the inner section through recursion
+ (output only). */
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
bool &maybe_zero_len, unsigned int &first_non_one,
bool &non_contiguous, enum c_omp_region_type ort,
- int *discontiguous)
+ int *discontiguous, tree *type)
{
- tree ret, low_bound, length, stride, type;
+ tree ret, low_bound, length, stride;
bool openacc = (ort & C_ORT_ACC) != 0;
if (TREE_CODE (t) != OMP_ARRAY_SECTION)
{
@@ -6113,6 +6114,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
if (type_dependent_expression_p (ret))
return NULL_TREE;
ret = convert_from_reference (ret);
+ *type = TREE_TYPE (ret);
return ret;
}
@@ -6124,14 +6126,10 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
TREE_OPERAND (t, 0) = omp_privatize_field (TREE_OPERAND (t, 0), false);
ret = handle_omp_array_sections_1 (c, TREE_OPERAND (t, 0), types,
maybe_zero_len, first_non_one,
- non_contiguous, ort, discontiguous);
+ non_contiguous, ort, discontiguous, type);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
- if (TREE_CODE (ret) == OMP_ARRAY_SECTION)
- type = TREE_TYPE (TREE_TYPE (TREE_OPERAND (ret, 0)));
- else
- type = TREE_TYPE (ret);
low_bound = TREE_OPERAND (t, 1);
length = TREE_OPERAND (t, 2);
stride = TREE_OPERAND (t, 3);
@@ -6235,11 +6233,11 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
&& (TREE_CODE (length) != INTEGER_CST || integer_onep (length)))
first_non_one++;
}
- if (TREE_CODE (type) == ARRAY_TYPE)
+ if (TREE_CODE (*type) == ARRAY_TYPE)
{
if (length == NULL_TREE
- && (TYPE_DOMAIN (type) == NULL_TREE
- || TYPE_MAX_VALUE (TYPE_DOMAIN (type)) == NULL_TREE))
+ && (TYPE_DOMAIN (*type) == NULL_TREE
+ || TYPE_MAX_VALUE (TYPE_DOMAIN (*type)) == NULL_TREE))
{
error_at (OMP_CLAUSE_LOCATION (c),
"for unknown bound array type length expression must "
@@ -6263,13 +6261,11 @@ 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 (TYPE_DOMAIN (type)
- && TYPE_MAX_VALUE (TYPE_DOMAIN (type))
- && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (type)))
- == INTEGER_CST)
+ if (TYPE_DOMAIN (*type) && TYPE_MAX_VALUE (TYPE_DOMAIN (*type))
+ && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (*type))) == INTEGER_CST)
{
tree size
- = fold_convert (sizetype, TYPE_MAX_VALUE (TYPE_DOMAIN (type)));
+ = fold_convert (sizetype, TYPE_MAX_VALUE (TYPE_DOMAIN (*type)));
size = size_binop (PLUS_EXPR, size, size_one_node);
if (TREE_CODE (low_bound) == INTEGER_CST)
{
@@ -6296,11 +6292,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
}
maybe_zero_len = true;
}
- else if (length == NULL_TREE
- && first_non_one == types.length ()
- && tree_int_cst_equal
- (TYPE_MAX_VALUE (TYPE_DOMAIN (type)),
- low_bound))
+ else if (length == NULL_TREE && first_non_one == types.length ()
+ && tree_int_cst_equal (
+ TYPE_MAX_VALUE (TYPE_DOMAIN (*type)), low_bound))
first_non_one++;
}
else if (length == NULL_TREE)
@@ -6382,7 +6376,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
}
}
}
- else if (TYPE_PTR_P (type))
+ else if (TYPE_PTR_P (*type))
{
if (length == NULL_TREE)
{
@@ -6459,7 +6453,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND)
- types.safe_push (type);
+ types.safe_push (*type);
/* We will need to evaluate lb more than once. */
tree lb = cp_save_expr (low_bound);
if (lb != low_bound)
@@ -6488,6 +6482,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
else
ret = grok_array_decl (OMP_CLAUSE_LOCATION (c), ret, low_bound, NULL,
tf_warning_or_error);
+
+ *type = TREE_TYPE (*type);
return ret;
}
@@ -6528,9 +6524,10 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort,
|| 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,
- non_contiguous, ort, discontiguous);
+ tree type;
+ tree first
+ = handle_omp_array_sections_1 (c, *tp, types, maybe_zero_len, first_non_one,
+ non_contiguous, ort, discontiguous, &type);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -6573,6 +6570,11 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort,
return false;
bool higher_discontiguous = false;
+ /* Whether the first dimension uses the extended array section syntax,
+ i.e. whether it specifies a length or a stride. */
+ bool first_extended = false;
+ /* Whether any dimension container except the last is a pointer. */
+ bool non_last_pointer = false;
for (i = num, t = OMP_CLAUSE_DECL (c); i > 0;
t = TREE_OPERAND (t, 0))
@@ -6638,10 +6640,15 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort,
full_span = true;
}
+ tree op = TREE_OPERAND (t, 0);
+ bool shaped
+ = TREE_CODE (op) == VIEW_CONVERT_EXPR
+ && TREE_CODE (TREE_TYPE (op)) == ARRAY_TYPE
+ && TREE_CODE (TREE_OPERAND (op, 0)) == OMP_ARRAY_SECTION;
if (!integer_onep (stride)
|| (higher_discontiguous
- && (!integer_zerop (low_bound)
- || !full_span)))
+ && (!integer_zerop (low_bound) || !full_span))
+ || shaped)
*discontiguous = 2;
if (!integer_onep (stride)
@@ -6652,20 +6659,29 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort,
if (!maybe_zero_len && i > first_non_one)
{
+ bool full_span;
if (integer_nonzerop (low_bound))
goto is_noncontiguous;
- if (length != NULL_TREE
- && TREE_CODE (length) == INTEGER_CST
- && TYPE_DOMAIN (types[i])
- && TYPE_MAX_VALUE (TYPE_DOMAIN (types[i]))
- && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])))
- == INTEGER_CST)
+ if (length != NULL_TREE && TREE_CODE (length) == INTEGER_CST)
{
- tree size;
- size = size_binop (PLUS_EXPR,
- TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])),
- size_one_node);
- if (!tree_int_cst_equal (length, size))
+ /* A pointer-typed dimension has no static bound to compare
+ against, so it can never be proven to span the whole
+ dimension; treat it conservatively as not full span. */
+ full_span = false;
+ if (TREE_CODE (types[i]) == ARRAY_TYPE
+ && TYPE_DOMAIN (types[i])
+ && TYPE_MAX_VALUE (TYPE_DOMAIN (types[i]))
+ && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])))
+ == INTEGER_CST)
+ {
+ tree size;
+ size
+ = size_binop (PLUS_EXPR,
+ TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])),
+ size_one_node);
+ full_span = tree_int_cst_equal (length, size);
+ }
+ if (!full_span)
{
is_noncontiguous:
if (discontiguous && *discontiguous)
@@ -6743,9 +6759,26 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort,
else
size = size_binop (MULT_EXPR, size, l);
}
+
+ if (i == 0
+ && ((length != NULL_TREE && !integer_onep (length))
+ || !integer_onep (stride)))
+ first_extended = true;
+ if (i < num - 1 && TREE_CODE (types[i]) == POINTER_TYPE
+ && /* accept decayed array */ !(
+ i == 0 && TREE_CODE (TREE_TYPE (types[i])) == ARRAY_TYPE))
+ non_last_pointer = true;
}
+
if (!processing_template_decl)
{
+ if (first_extended && non_last_pointer)
+ error_at (
+ OMP_CLAUSE_LOCATION (c),
+ "only the last dimension can refer to a pointer when the first "
+ "dimension uses the extended array section syntax in %qs clause",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+
if (non_contiguous)
{
int kind = OMP_CLAUSE_MAP_KIND (c);
@@ -10329,7 +10362,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
}
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE)
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_SHAPE_DIM)
break;
/* FALLTHRU */
case OMP_CLAUSE_TO:
diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc
index 5dd594ac211..ab723c1202c 100644
--- a/gcc/omp-general.cc
+++ b/gcc/omp-general.cc
@@ -4229,8 +4229,9 @@ omp_parse_noncontiguous_array (tree *expr0)
tree expr = *expr0;
bool noncontig = false;
- while (TREE_CODE (expr) == OMP_ARRAY_SECTION
- || TREE_CODE (expr) == ARRAY_REF)
+ while (TREE_CODE (expr) == OMP_ARRAY_SECTION || TREE_CODE (expr) == ARRAY_REF
+ || (TREE_CODE (expr) == VIEW_CONVERT_EXPR
+ && TREE_CODE (TREE_OPERAND (expr, 0)) == OMP_ARRAY_SECTION))
{
/* Contiguous arrays use ARRAY_REF. By the time we reach here,
OMP_ARRAY_SECTION is only used for noncontiguous arrays. */
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..99756bbf4c2
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-1.c
@@ -0,0 +1,33 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */
+
+/* Test target update with a single level of pointer indirection, one dimension
+ before and after. */
+
+#define DIM1 5
+#define DIM2 10
+
+void fixed_index (void)
+{
+ int *x[DIM1];
+#pragma omp target update to(x[2][ :DIM2])
+/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 10\] \[pointer\]\)} "original" } } */
+}
+
+/* Index, length and stride need not be literal constants -- a variable
+ works just as well. */
+
+void range_of_pointers (int lo, int len)
+{
+ int *x[DIM1];
+#pragma omp target update to(x[lo:len][ :DIM2])
+/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:SAVE_EXPR <lo> \[len: len\]\) map\(grid_dim:0 \[len: 10\] \[pointer\]\)} "original" } } */
+}
+
+void strided_range_of_pointers (int lo, int len, int str)
+{
+ int *x[DIM1];
+#pragma omp target update to(x[lo:len:str][ :DIM2])
+/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:SAVE_EXPR <lo> \[len: len\]\) map\(grid_stride:str\) map\(grid_dim:0 \[len: 10\] \[pointer\]\)} "original" } } */
+}
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..15bc2cb28ba
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-2.c
@@ -0,0 +1,15 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */
+
+/* Test target update with a 2D array of pointers */
+
+#define DIM1 5
+#define DIM2 5
+
+void two_d_array_of_pointers (void)
+{
+ int *y[DIM1][DIM2];
+#pragma omp target update to(y[0][ :2])
+/* { dg-final { scan-tree-dump {map\(to_grid:y \[len: [0-9]+\]\) map\(grid_dim:0 \[len: 1\]\) map\(grid_dim:0 \[len: 2\]\)} "original" } } */
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-3.c b/gcc/testsuite/c-c++-common/gomp/array-section-3.c
new file mode 100644
index 00000000000..f7b17475b5d
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-3.c
@@ -0,0 +1,38 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */
+
+/* A 2D array-shaping cast applied to a 2D array of pointers -- two segments:
+ the array-of-pointers side, selected by two fixed indices, and the reshaped
+ pointee side, both of its dimensions explicitly sectioned. */
+
+#define DIM1 3
+#define DIM2 4
+#define ROWS 6
+#define COLS 6
+
+void shape_cast_two_fixed_indices_strided_row (void)
+{
+ int *x[DIM1][DIM2];
+#pragma omp target update to((([ROWS][COLS]) x[1][1])[1:3:2][0:2])
+/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 2\]\) map\(shape_dim:6\)} "original" } } */
+}
+
+void shape_cast_length_one_range_index (void)
+{
+ int *x[DIM1][DIM2];
+#pragma omp target update to((([ROWS][COLS]) x[2][0:1])[2:2][1:3])
+/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 1\]\) map\(grid_dim:2 \[len: 2\] \[pointer\]\) map\(shape_dim:6\) map\(grid_dim:1 \[len: 3\]\) map\(shape_dim:6\)} "original" } } */
+}
+
+/* Same shape as case 1, but with variables everywhere a literal is
+ allowed: the two fixed array-of-pointers indices, and the reshaped
+ pointee side's index/length/stride. */
+
+void shape_cast_var (int i, int j, int r0, int rlen, int rstr, int c0,
+ int clen)
+{
+ int *x[DIM1][DIM2];
+#pragma omp target update to((([ROWS][COLS]) x[i][j])[r0:rlen:rstr][c0:clen])
+/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:i \[len: 1\]\) map\(grid_dim:j \[len: 1\]\) map\(grid_dim:SAVE_EXPR <r0> \[len: rlen\] \[pointer\]\) map\(grid_stride:rstr\) map\(shape_dim:6\) map\(grid_dim:SAVE_EXPR <c0> \[len: clen\]\) map\(shape_dim:6\)} "original" } } */
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-4.c b/gcc/testsuite/c-c++-common/gomp/array-section-4.c
new file mode 100644
index 00000000000..ce446b931c2
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-4.c
@@ -0,0 +1,21 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */
+
+/* An array-of-pointers with several dimensions selected by consecutive fixed
+ indices before the final range.
+ Confirms walking several plain ARRAY_REFs (rather than OMP_ARRAY_SECTIONs)
+ before reaching the pointer dereference produces one GOMP_MAP_GRID_DIM per
+ fixed index, plus a final one for the pointer crossing. */
+
+#define D1 3
+#define D2 3
+#define D3 4
+#define M 8
+
+void three_consecutive_fixed_indices (void)
+{
+ int *w[D1][D2][D3];
+#pragma omp target update to(w[1][2][0][2:4])
+/* { dg-final { scan-tree-dump {map\(to_grid:w \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 1\]\) map\(grid_dim:2 \[len: 4\] \[pointer\]\)} "original" } } */
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-5.c b/gcc/testsuite/c-c++-common/gomp/array-section-5.c
new file mode 100644
index 00000000000..9bd8a251393
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-5.c
@@ -0,0 +1,25 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */
+
+/* Test an array-shaping cast that has more dimensions than are explicitly
+ sectioned afterwards -- copy the whole span gathered from the cast. */
+
+#define DIM1 3
+#define DIM2 4
+#define ROWS 6
+#define COLS 6
+
+void shape_cast_omitted_column_dim_strided (void)
+{
+ int *x[DIM1][DIM2];
+#pragma omp target update to((([ROWS][COLS]) x[1][1])[1:3:2])
+/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 6\]\) map\(shape_dim:6\)} "original" } } */
+}
+
+void shape_cast_omitted_column_dim_length_one_range (void)
+{
+ int *x[DIM1][DIM2];
+#pragma omp target update to((([ROWS][COLS]) x[2][0:1])[2:2])
+/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 1\]\) map\(grid_dim:2 \[len: 2\] \[pointer\]\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 6\]\) map\(shape_dim:6\)} "original" } } */
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-6.c b/gcc/testsuite/c-c++-common/gomp/array-section-6.c
new file mode 100644
index 00000000000..e90f4e66188
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-6.c
@@ -0,0 +1,18 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */
+
+/* An array-shaping cast applied to a plain pointer (no array-of-pointers
+ indirection involved), with a further, unit-stride, contiguous section
+ applied outside the cast. Check that it simplifies correctly and does not
+ go through the noncontiguous-update codepath. */
+
+#define N 100
+
+void plain_pointer_shape_cast (void)
+{
+ int *ptr;
+#pragma omp target update to((([N]) ptr)[10:30])
+/* { dg-final { scan-tree-dump {to\(VIEW_CONVERT_EXPR<int\[100\]>\(\*ptr\)\[10\] \[len: 120\]\)} "original" } } */
+/* { dg-final { scan-tree-dump-not "to_grid" "original" } } */
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-8.c b/gcc/testsuite/c-c++-common/gomp/array-section-8.c
new file mode 100644
index 00000000000..4873138362a
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-8.c
@@ -0,0 +1,27 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */
+
+/* Check three-segment (two pointer indirections) noncontiguous update: a
+ strided section reached by crossing two levels of pointer indirection. */
+
+#define DIM1 4
+
+void three_segment_one_dim_per_segment (void)
+{
+ int **arr[DIM1];
+#pragma omp target update from(arr[2][0:2:2][3:5:2])
+/* { dg-final { scan-tree-dump {map\(from_grid:arr \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 2\] \[pointer\]\) map\(grid_stride:2\) map\(grid_dim:3 \[len: 5\] \[pointer\]\) map\(grid_stride:2\)} "original" } } */
+}
+
+/* Same shape, but every index/length/stride across all three segments is
+ a variable rather than a literal. */
+
+void three_segment_one_dim_per_segment_var (int p, int p0, int plen,
+ int pstr, int e0, int elen,
+ int estr)
+{
+ int **arr[DIM1];
+#pragma omp target update from(arr[p][p0:plen:pstr][e0:elen:estr])
+/* { dg-final { scan-tree-dump {map\(from_grid:arr \[len: [0-9]+\]\) map\(grid_dim:SAVE_EXPR <p> \[len: 1\]\) map\(grid_dim:SAVE_EXPR <p0> \[len: plen\] \[pointer\]\) map\(grid_stride:pstr\) map\(grid_dim:SAVE_EXPR <e0> \[len: elen\] \[pointer\]\) map\(grid_stride:estr\)} "original" } } */
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-9.c b/gcc/testsuite/c-c++-common/gomp/array-section-9.c
new file mode 100644
index 00000000000..9d2e0ceae09
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-9.c
@@ -0,0 +1,21 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-lower" } */
+
+/* Check three-segment (two pointer indirections) noncontiguous update with two
+ dimensions selected per segment -- six dimensions total. */
+
+#define DIM1A 2
+#define DIM1B 3
+#define ROWLEN 20
+#define NCOLS 3
+
+typedef int (*p2_t)[ROWLEN];
+typedef p2_t (*p1_t)[NCOLS];
+
+void three_segment_two_dims_per_segment (void)
+{
+ p1_t arr[DIM1A][DIM1B];
+#pragma omp target update from(arr[1][0:2:2][0:2:2][0:2:2][0:2][3:5:2])
+/* { dg-final { scan-tree-dump {map\(from_grid:arr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:0 \[len: 2\]\) map\(grid_stride:2\) map\(grid_dim:0 \[len: 2\] \[pointer\]\) map\(grid_stride:2\) map\(grid_dim:0 \[len: 2\]\) map\(grid_stride:2\) map\(grid_dim:0 \[len: 2\] \[pointer\]\) map\(grid_dim:3 \[len: 5\]\) map\(grid_stride:2\)} "original" } } */
+}
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
new file mode 100644
index 00000000000..fa2fccd93c6
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c
@@ -0,0 +1,19 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+/* Regression test for strided target updates with iterators. */
+
+#define DIM1 17
+#define DIM2 39
+
+void f (int *x[DIM1])
+{
+ /* Per-index iterator selecting the pointer, no stride on the pointee:
+ still supported. */
+#pragma omp target update to (iterator(i=0:DIM1): x[i][ :DIM2])
+
+ /* Per-index iterator selecting the pointer, but an explicit stride on
+ the pointee: still sorry. */
+#pragma omp target update to (iterator(i=0:DIM1): x[i][0:DIM2:2])
+ /* { dg-message "sorry, unimplemented: strided target updates with iterators" "" { target *-*-* } .-1 } */
+}
diff --git a/gcc/testsuite/g++.dg/gomp/array-shaping-3.C b/gcc/testsuite/g++.dg/gomp/array-shaping-3.C
new file mode 100644
index 00000000000..423b298d4ff
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/array-shaping-3.C
@@ -0,0 +1,71 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-original" }
+
+/* Test target update on array-shaping cast applied to a section selected from a
+ 2D array-of-pointers base, when the enclosing function is a template. */
+
+#define DIM1 7
+#define DIM2 9
+#define ROWS 6
+#define COLS 6
+
+template<typename T>
+void foo()
+{
+ T *x[DIM1][DIM2];
+
+#pragma omp target update to((([ROWS][COLS]) x[1][2])[1:3:2][0:2])
+// { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 2\]\) map\(shape_dim:6\)} "original" } }
+// { dg-final { scan-tree-dump-not {grid_dim:0 \[len: 9\]} "original" } }
+// { dg-final { scan-tree-dump-not {shape_dim:9} "original" } }
+}
+
+template<typename T>
+void bar()
+{
+ T *x[DIM1];
+
+#pragma omp target update to((([ROWS][COLS]) x[1])[1:3:2][0:2])
+// { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 2\]\) map\(shape_dim:6\)} "original" } }
+// { dg-final { scan-tree-dump-not {grid_dim:0 \[len: 9\]} "original" } }
+// { dg-final { scan-tree-dump-not {shape_dim:9} "original" } }
+}
+
+void control()
+{
+ int *x[DIM1][DIM2];
+
+#pragma omp target update to((([ROWS][COLS]) x[1][2])[1:3:2][0:2])
+// { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 2\]\) map\(shape_dim:6\)} "original" } }
+}
+
+/* Reference + template combination: same 2D array-of-pointers
+ base as foo() above, but accessed through a C++ reference. */
+
+template<typename T>
+void baz()
+{
+ T *x[DIM1][DIM2];
+ T *(&x_ref)[DIM1][DIM2] = x;
+
+#pragma omp target update to((([ROWS][COLS]) x_ref[1][2])[1:3:2][0:2])
+// { dg-final { scan-tree-dump {map\(to_grid:\*x_ref \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 2\]\) map\(shape_dim:6\)} "original" } }
+// { dg-final { scan-tree-dump-not {grid_dim:0 \[len: 9\]} "original" } }
+// { dg-final { scan-tree-dump-not {shape_dim:9} "original" } }
+}
+
+void control_ref()
+{
+ int *x[DIM1][DIM2];
+ int *(&x_ref)[DIM1][DIM2] = x;
+
+#pragma omp target update to((([ROWS][COLS]) x_ref[1][2])[1:3:2][0:2])
+// { dg-final { scan-tree-dump {map\(to_grid:\*x_ref \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 1\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:6\) map\(grid_dim:0 \[len: 2\]\) map\(shape_dim:6\)} "original" } }
+}
+
+void g()
+{
+ foo<int> ();
+ bar<int> ();
+ baz<int> ();
+}
diff --git a/gcc/testsuite/g++.dg/gomp/array-shaping-4.C b/gcc/testsuite/g++.dg/gomp/array-shaping-4.C
new file mode 100644
index 00000000000..e857ff01c1e
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/array-shaping-4.C
@@ -0,0 +1,37 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-original" }
+
+/* Test an array-shaping cast applied directly to a single element of an
+ array of pointers, with and without outer array section.
+ Also check templates and variable subscript indices. */
+
+#define DIM1 5
+#define N 20
+
+template<typename T>
+void foo()
+{
+ T *x[DIM1];
+
+#pragma omp target update to(([N]) x[2])
+// { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 20\] \[pointer\]\) map\(shape_dim:20\)} "original" } }
+
+#pragma omp target update to((([N][N+1]) x[2])[1:3:2][1:4])
+// { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:20\) map\(grid_dim:1 \[len: 4\]\) map\(shape_dim:21\)} "original" } }
+}
+
+void control()
+{
+ int *x[DIM1];
+
+#pragma omp target update to(([N]) x[2])
+// { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 20\] \[pointer\]\) map\(shape_dim:20\)} "original" } }
+
+#pragma omp target update to((([N][N+1]) x[2])[1:3:2][1:4])
+// { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:20\) map\(grid_dim:1 \[len: 4\]\) map\(shape_dim:21\)} "original" } }
+}
+
+void g()
+{
+ foo<int> ();
+}
diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-9.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-9.C
new file mode 100644
index 00000000000..c8ce6f61fe0
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-9.C
@@ -0,0 +1,51 @@
+// { dg-do compile }
+
+/* Test rejected array-shaping cast combinations. */
+
+#define DIM1 5
+#define DIM2 5
+
+template<typename T>
+void foo ()
+{
+ T *x[DIM1];
+
+#pragma omp target update to(([DIM2]) x[1:3][0:5])
+// { dg-error "OpenMP array shaping operator with non-pointer argument" "" { target *-*-* } .-1 }
+// { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 }
+
+ T **y[DIM1];
+
+#pragma omp target update to(([DIM2]) y[0:2])
+// { dg-error "OpenMP array shaping operator with non-pointer argument" "" { target *-*-* } .-1 }
+// { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 }
+
+ /* A *fixed* index into the same array-of-pointers-of-pointers -- unlike the range above, this is not rejected by the non-pointer-argument check: a fixed index does denote a single pointer expression. But reshaping it here would still leave a pointer-typed result ("T *[DIM2]"): copying that host-to-device as plain data would copy raw host pointer values with no attach translation, which is unsupported (for now). */
+#pragma omp target update to(([DIM2]) y[2])
+ // { dg-message "sorry, unimplemented: more than one level of pointer indirection in a noncontiguous target update" "" { target *-*-* } .-1 }
+ // { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 }
+}
+
+void control ()
+{
+ int *x[DIM1];
+
+#pragma omp target update to(([DIM2]) x[1:3][0:5])
+// { dg-error "OpenMP array shaping operator with non-pointer argument" "" { target *-*-* } .-1 }
+// { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 }
+
+ int **y[DIM1];
+
+#pragma omp target update to(([DIM2]) y[0:2])
+// { dg-error "OpenMP array shaping operator with non-pointer argument" "" { target *-*-* } .-1 }
+// { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 }
+
+#pragma omp target update to(([DIM2]) y[2])
+// { dg-message "sorry, unimplemented: more than one level of pointer indirection in a noncontiguous target update" "" { target *-*-* } .-1 }
+// { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 }
+}
+
+void g ()
+{
+ foo<int> ();
+}
diff --git a/gcc/testsuite/gcc.dg/gomp/array-shaping-9.c b/gcc/testsuite/gcc.dg/gomp/array-shaping-9.c
new file mode 100644
index 00000000000..9050a56515a
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/array-shaping-9.c
@@ -0,0 +1,25 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+
+#define DIM1 5
+#define N 20
+
+int main ()
+{
+ int *x[DIM1];
+
+ /* An array-shaping cast applied directly to a single element of an
+ array of pointers, with no further outer section. */
+
+#pragma omp target update to(([N]) x[2])
+/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:0 \[len: 20\] \[pointer\]\) map\(shape_dim:20\)} "original" } } */
+
+ /* Same base, but with an explicit outer section applied to the
+ reshaped 2D result. */
+
+#pragma omp target update to((([N][N+1]) x[2])[1:3:2][1:4])
+/* { dg-final { scan-tree-dump {map\(to_grid:x \[len: [0-9]+\]\) map\(grid_dim:2 \[len: 1\]\) map\(grid_dim:1 \[len: 3\] \[pointer\]\) map\(grid_stride:2\) map\(shape_dim:20\) map\(grid_dim:1 \[len: 4\]\) map\(shape_dim:21\)} "original" } } */
+
+ return 0;
+}
diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-section-c-9.c b/gcc/testsuite/gcc.dg/gomp/bad-array-section-c-9.c
new file mode 100644
index 00000000000..e2bf421baf8
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/bad-array-section-c-9.c
@@ -0,0 +1,60 @@
+/* { dg-do compile } */
+
+/* Test for bad array section syntax in OMP target update clauses implying
+ pointer indirection. */
+
+#include <stdlib.h>
+
+#define DIM1 5
+#define DIM2 5
+
+int main (void)
+{
+ int **z;
+
+ z = malloc(DIM1 * sizeof z);
+ for (int i = 0; i < DIM1; i++)
+ z[i] = malloc(DIM2 * sizeof(int));
+
+ #pragma omp target update to(z[:DIM1][0])
+ /* { dg-error "only the last dimension can refer to a pointer when the first dimension uses the extended array section syntax" "" { target *-*-* } .-1 } */
+
+ return 0;
+}
+
+#define D1 3
+#define D2 4
+#define ROWS 6
+#define COLS 6
+
+void
+g (void)
+{
+ int *x[D1][D2];
+ int i, j;
+
+ for (i = 0; i < D1; i++)
+ for (j = 0; j < D2; j++)
+ x[i][j] = malloc (ROWS * COLS * sizeof (int));
+
+ #pragma omp target update to(x[1][1:2][1:3][0:2])
+ /* { dg-error "does not have pointer or array type" "" { target *-*-* } .-1 } */
+ /* { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } */
+}
+
+void
+h (void)
+{
+ int **y[D1];
+ int i, j;
+
+ for (i = 0; i < D1; i++)
+ {
+ y[i] = malloc (D2 * sizeof (int *));
+ for (j = 0; j < D2; j++)
+ y[i][j] = malloc (10 * sizeof (int));
+ }
+
+ #pragma omp target update to(y[0:2][0:3][0:5])
+ /* { dg-error "only the last dimension can refer to a pointer when the first dimension uses the extended array section syntax" "" { target *-*-* } .-1 } */
+}
diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-8.c b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-8.c
new file mode 100644
index 00000000000..b05b182f2f5
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-8.c
@@ -0,0 +1,45 @@
+/* { dg-do compile } */
+
+/* Test rejected array-shaping cast combinations. */
+
+#define DIM1 5
+#define DIM2 5
+
+int main (void)
+{
+ int *x[DIM1];
+
+ /* Array-shaping operator applied on an array section. */
+#pragma omp target update to(([DIM2]) x[1:3][0:5])
+ /* { dg-error "OpenMP array shaping operator with non-pointer argument" "" { target *-*-* } .-1 } */
+
+ return 0;
+}
+
+/* Same restriction, on an array-of-pointers-of-pointers. */
+
+void
+depth2 (void)
+{
+ int **y[DIM1];
+
+#pragma omp target update to(([DIM2]) y[0:2])
+/* { dg-error "OpenMP array shaping operator with non-pointer argument" "" { target *-*-* } .-1 } */
+}
+
+/* A *fixed* index into the same array-of-pointers-of-pointers, e.g.
+ "y[2]" -- unlike the range above, this is not rejected by the
+ non-pointer-argument check. But reshaping it here would still leave
+ a pointer-typed result ("int *[DIM2]"): copying that host-to-device
+ as plain data would copy raw host pointer values with no attach
+ translation. */
+
+void
+depth2_fixed_index (void)
+{
+ int **y[DIM1];
+
+#pragma omp target update to(([DIM2]) y[2])
+ /* { dg-message "sorry, unimplemented: more than one level of pointer indirection in a noncontiguous target update" "" { target *-*-* } .-1 } */
+ /* { dg-error "must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } */
+}
--
2.53.0