Extend the grid descriptor lowering to build the noncontiguous-array
descriptor for array-shaping casts and array-of-pointers sections, not
just plain array sections:
- Support multiple segments, each starting where a GRID_DIM clause is
marked as selecting through a pointer (OMP_CLAUSE_MAP_GRID_DIM_POINTER),
tracking per-segment dimension/pointer counts (__nsegments,
__seg_ndims, __seg_nptrs) in the runtime descriptor.
- Consume GOMP_MAP_SHAPE_DIM clauses to size dimensions introduced by an
array-shaping cast, once the decl's own array type runs out of
dimensions.
- Stop walking GRID_DIM/GRID_STRIDE/SHAPE_DIM clauses once the last one
for the current segment has been consumed, rather than assuming a
fixed dimension count.
- Accept decayed array parameters (no array type left to consult) by
falling back to the last pointer type crossed.
- Only mark the runtime descriptor fields TREE_STATIC when every
contributing dimension is truly constant, instead of unconditionally.
Teach the clause-splicing logic in gimplify.cc about the new
GOMP_MAP_SHAPE_DIM clause, and add the "lower" dump scans checking the
generated descriptors for the tests added alongside the front-end
support.
gcc/ChangeLog:
* gimplify.cc (omp_group_last): Handle GOMP_MAP_SHAPE_DIM.
(gimplify_adjust_omp_clauses): Likewise.
* omp-low.cc (omp_noncontig_descriptor_type): Add __nsegments,
__seg_ndims and __seg_nptrs fields.
(lower_omp_target_grid_desc): Support multiple segments, dimensions
introduced by an array-shaping cast, decayed array parameters, and
only mark the descriptor fields static when every dimension is
constant.
(lower_omp_target): Handle GOMP_MAP_SHAPE_DIM alongside
GOMP_MAP_GRID_DIM/GOMP_MAP_GRID_STRIDE.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/array-section-1.c: Add "lower" dump scans.
* c-c++-common/gomp/array-section-2.c: Likewise.
* c-c++-common/gomp/array-section-3.c: Likewise.
* c-c++-common/gomp/array-section-4.c: Likewise.
* c-c++-common/gomp/array-section-5.c: Likewise.
* c-c++-common/gomp/array-section-6.c: Likewise.
* c-c++-common/gomp/array-section-8.c: Likewise.
* c-c++-common/gomp/array-section-9.c: Likewise.
---
gcc/gimplify.cc | 18 +-
gcc/omp-low.cc | 462 ++++++++++++------
.../c-c++-common/gomp/array-section-1.c | 7 +
.../c-c++-common/gomp/array-section-2.c | 7 +
.../c-c++-common/gomp/array-section-3.c | 7 +
.../c-c++-common/gomp/array-section-4.c | 7 +
.../c-c++-common/gomp/array-section-5.c | 7 +
.../c-c++-common/gomp/array-section-6.c | 1 +
.../c-c++-common/gomp/array-section-8.c | 7 +
.../c-c++-common/gomp/array-section-9.c | 7 +
10 files changed, 361 insertions(+), 169 deletions(-)
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index b0ecf08cf7f..e4671befc55 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -11608,10 +11608,10 @@ omp_group_last (tree *start_p)
case GOMP_MAP_TO_GRID:
case GOMP_MAP_FROM_GRID:
- while (nc
- && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+ while (nc && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM
- || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE))
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_SHAPE_DIM))
{
grp_last_p = &OMP_CLAUSE_CHAIN (c);
c = nc;
@@ -16933,7 +16933,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
break;
if (OMP_CLAUSE_SIZE (c) == NULL_TREE
&& 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)
{
/* Sanity check: attach/detach map kinds use the size as a bias,
and it's never right to use the decl size for such
@@ -17036,11 +17037,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
remove = true;
}
else 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)
{
- /* The OMP_CLAUSE_DECL for GRID_DIM/GRID_STRIDE isn't necessarily
- an lvalue -- e.g. it might be a constant. So handle it
- specially here. */
+ /* The OMP_CLAUSE_DECL for GRID_DIM/GRID_STRIDE/SHAPE_DIM isn't
+ necessarily an lvalue -- e.g. it might be a constant. So
+ handle it specially here. */
if (gimplify_expr (&OMP_CLAUSE_DECL (c), seq_p, NULL,
is_gimple_val, fb_rvalue) == GS_ERROR)
{
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 45f6c701f48..128a0848ad3 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1355,6 +1355,21 @@ omp_noncontig_descriptor_type (location_t loc)
TREE_CHAIN (field) = fields;
fields = field;
+ field = build_decl (loc, FIELD_DECL, get_identifier ("__nsegments"),
+ size_type_node);
+ TREE_CHAIN (field) = fields;
+ fields = field;
+
+ field = build_decl (loc, FIELD_DECL, get_identifier ("__seg_ndims"),
+ ptr_size_type);
+ TREE_CHAIN (field) = fields;
+ fields = field;
+
+ field = build_decl (loc, FIELD_DECL, get_identifier ("__seg_nptrs"),
+ ptr_size_type);
+ TREE_CHAIN (field) = fields;
+ fields = field;
+
finish_builtin_struct (t, "__omp_noncontig_desc_type", fields,
ptr_type_node);
cached = t;
@@ -13723,52 +13738,111 @@ convert_from_firstprivate_int (tree var, tree
orig_type, bool is_ref,
gimplify_assign (tmp, var, gs);
return fold_build1 (VIEW_CONVERT_EXPR, type, tmp);
+}
/* Build the noncontiguous-array descriptor for the
- GOMP_MAP_TO_GRID/GOMP_MAP_FROM_GRID clause C (whose GOMP_MAP_TO_PSET
+ GOMP_MAP_TO_GRID/GOMP_MAP_FROM_GRID clause C. The GOMP_MAP_TO_PSET
sibling clause holds the descriptor decl and the GOMP_MAP_GRID_DIM/
- GOMP_MAP_GRID_STRIDE clauses that follow describe each dimension). Any
- gimplification needed to initialize the descriptor is appended to
+ GOMP_MAP_GRID_STRIDE/GOMP_MAP_SHAPE_DIM clauses that follow describe each
+ dimension. Then splice those consumed dimension clauses out of the clause
+ chain. Any gimplification needed to initialize the descriptor is appended
to
*ILIST_P. */
static void
lower_omp_target_grid_desc (tree c, gimple_seq *ilist_p)
{
tree decl = OMP_CLAUSE_DECL (c);
- tree dn = OMP_CLAUSE_CHAIN (c);
- gcc_assert (OMP_CLAUSE_CODE (dn) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (dn) == GOMP_MAP_TO_PSET);
- tree desc = OMP_CLAUSE_DECL (dn);
-
+ tree desc_node = OMP_CLAUSE_CHAIN (c);
+ gcc_assert (OMP_CLAUSE_CODE (desc_node) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (desc_node) == GOMP_MAP_TO_PSET);
+ tree desc = OMP_CLAUSE_DECL (desc_node);
tree oc, elsize = OMP_CLAUSE_SIZE (c);
tree type = TREE_TYPE (decl);
- int i, dims = 0;
+
+ /* First, count dimensions and record their types. */
+
+ int i, dims = 0, segments = 0;
tree nc;
auto_vec<tree> tdims;
+ auto_vec<int> seg_ndims, seg_nptrs;
bool pointer_based = false, handled_pointer_section = false;
- tree arrsize = size_one_node;
- /* Allow a single (maybe strided) array section if we have a
- pointer base. */
if (TREE_CODE (decl) == INDIRECT_REF
- && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
- == POINTER_TYPE))
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) == POINTER_TYPE))
{
+ /* Allow a single (maybe strided) array section if we have a
+ pointer base. */
pointer_based = true;
dims = 1;
+ seg_ndims.safe_push (dims);
+ segments = 1;
}
else
- /* NOTE: Don't treat (e.g. Fortran, fixed-length) strings as
- array types here; array section syntax isn't applicable to
- strings. */
- for (tree itype = type;
- TREE_CODE (itype) == ARRAY_TYPE
- && !TYPE_STRING_FLAG (itype);
- itype = TREE_TYPE (itype))
- {
- tdims.safe_push (itype);
- dims++;
- }
+ {
+ tree dim_clause = OMP_CLAUSE_CHAIN (desc_node);
+ gcc_assert (OMP_CLAUSE_CODE (dim_clause) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (dim_clause) == GOMP_MAP_GRID_DIM);
+ int seg_dims = 0;
+
+ tree ptr_type = NULL_TREE;
+ tree itype = type;
+ bool have_itype;
+ while ((have_itype = (TREE_CODE (itype) == ARRAY_TYPE &&
+ /*array section syntax isn't applicable to strings*/
+ !TYPE_STRING_FLAG (itype))
+ || TREE_CODE (itype) == POINTER_TYPE)
+ || dim_clause)
+ {
+ /* Once the decl type runs out of dimensions
+ (e.g. with an array-shaping cast), treat further
+ dim_clauses as having the type of the last pointer
+ crossed. */
+ tree cur_type = have_itype ? itype : ptr_type;
+ gcc_assert (cur_type);
+
+ if (dim_clause == NULL_TREE && TREE_CODE (cur_type) == POINTER_TYPE)
+ {
+ elsize = TYPE_SIZE_UNIT (cur_type);
+ break;
+ }
+ tdims.safe_push (cur_type);
+ if (TREE_CODE (cur_type) == POINTER_TYPE)
+ ptr_type = cur_type;
+
+ if (dim_clause && OMP_CLAUSE_MAP_GRID_DIM_POINTER (dim_clause))
+ {
+ gcc_assert (POINTER_TYPE_P (cur_type));
+ seg_ndims.safe_push (seg_dims);
+ seg_dims = 1;
+ segments++;
+ }
+ else
+ seg_dims++;
+ dims++;
+
+ if (dim_clause)
+ {
+ dim_clause = OMP_CLAUSE_CHAIN (dim_clause);
+ while (
+ dim_clause
+ && (OMP_CLAUSE_CODE (dim_clause) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (dim_clause) == GOMP_MAP_GRID_STRIDE
+ || OMP_CLAUSE_MAP_KIND (dim_clause)
+ == GOMP_MAP_SHAPE_DIM)))
+ dim_clause = OMP_CLAUSE_CHAIN (dim_clause);
+
+ if (dim_clause
+ && (OMP_CLAUSE_CODE (dim_clause) != OMP_CLAUSE_MAP
+ || OMP_CLAUSE_MAP_KIND (dim_clause) != GOMP_MAP_GRID_DIM))
+ dim_clause = NULL_TREE;
+ }
+
+ if (have_itype)
+ itype = TREE_TYPE (itype);
+ }
+ seg_ndims.safe_push (seg_dims);
+ segments++;
+ }
unsigned tdim = 0;
@@ -13776,156 +13850,199 @@ lower_omp_target_grid_desc (tree c, gimple_seq
*ilist_p)
vec<constructor_elt, va_gc> *vindex;
vec<constructor_elt, va_gc> *vlen;
vec<constructor_elt, va_gc> *vstride;
+ vec<constructor_elt, va_gc> *vseg_ndims;
+ vec<constructor_elt, va_gc> *vseg_nptrs;
vec_alloc (vdim, dims);
vec_alloc (vindex, dims);
vec_alloc (vlen, dims);
vec_alloc (vstride, dims);
+ vec_alloc (vseg_ndims, segments);
+ vec_alloc (vseg_nptrs, segments - 1);
- tree size_arr_type
- = build_array_type_nelts (size_type_node, dims);
+ tree dim_arr_type = build_array_type_nelts (size_type_node, dims);
+ tree seg_arr_type = build_array_type_nelts (size_type_node, segments);
+ tree seg1_arr_type
+ = build_array_type_nelts (size_type_node, MAX (segments - 1, 1));
- tree dim_tmp = create_tmp_var (size_arr_type, ".omp_dim");
+ tree dim_tmp = create_tmp_var (dim_arr_type, ".omp_dim");
DECL_NAMELESS (dim_tmp) = 1;
TREE_ADDRESSABLE (dim_tmp) = 1;
TREE_STATIC (dim_tmp) = 1;
- tree index_tmp = create_tmp_var (size_arr_type, ".omp_index");
+ tree index_tmp = create_tmp_var (dim_arr_type, ".omp_index");
DECL_NAMELESS (index_tmp) = 1;
TREE_ADDRESSABLE (index_tmp) = 1;
TREE_STATIC (index_tmp) = 1;
- tree len_tmp = create_tmp_var (size_arr_type, ".omp_len");
+ tree len_tmp = create_tmp_var (dim_arr_type, ".omp_len");
DECL_NAMELESS (len_tmp) = 1;
TREE_ADDRESSABLE (len_tmp) = 1;
TREE_STATIC (len_tmp) = 1;
- tree stride_tmp = create_tmp_var (size_arr_type, ".omp_stride");
+ tree stride_tmp = create_tmp_var (dim_arr_type, ".omp_stride");
DECL_NAMELESS (stride_tmp) = 1;
TREE_ADDRESSABLE (stride_tmp) = 1;
TREE_STATIC (stride_tmp) = 1;
+ tree seg_ndims_tmp = create_tmp_var (seg_arr_type, ".omp_seg_ndims");
+ DECL_NAMELESS (seg_ndims_tmp) = 1;
+ TREE_ADDRESSABLE (seg_ndims_tmp) = 1;
+ TREE_STATIC (seg_ndims_tmp) = 1;
+ tree seg_nptrs_tmp = create_tmp_var (seg1_arr_type, ".omp_seg_nptrs");
+ DECL_NAMELESS (seg_nptrs_tmp) = 1;
+ TREE_ADDRESSABLE (seg_nptrs_tmp) = 1;
+ TREE_STATIC (seg_nptrs_tmp) = 1;
oc = c;
- c = dn;
+ c = desc_node;
tree span = NULL_TREE;
- for (i = 0; i < dims; i++)
+ for (int j = 0, idim = 0; j < segments; j++)
{
- nc = OMP_CLAUSE_CHAIN (c);
- tree dim = NULL_TREE, index = NULL_TREE, len = NULL_TREE,
- stride = size_one_node;
+ tree nptrs = size_int (1);
- if (nc
- && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM)
+ for (i = 0; i < seg_ndims[j]; i++, idim++)
{
- index = OMP_CLAUSE_DECL (nc);
- len = OMP_CLAUSE_SIZE (nc);
+ nc = OMP_CLAUSE_CHAIN (c);
+ tree dim = NULL_TREE, index = NULL_TREE, len = NULL_TREE,
+ stride = size_one_node;
- index = fold_convert (sizetype, index);
- len = fold_convert (sizetype, len);
-
- tree nc2 = OMP_CLAUSE_CHAIN (nc);
- if (nc2
- && OMP_CLAUSE_CODE (nc2) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (nc2)
- == GOMP_MAP_GRID_STRIDE))
+ if (nc && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM)
{
- stride = OMP_CLAUSE_DECL (nc2);
- stride = fold_convert (sizetype, stride);
- if (OMP_CLAUSE_SIZE (nc2))
+ index = OMP_CLAUSE_DECL (nc);
+ len = OMP_CLAUSE_SIZE (nc);
+
+ index = fold_convert (sizetype, index);
+ len = fold_convert (sizetype, len);
+
+ nptrs = size_binop (MULT_EXPR, nptrs, len);
+
+ tree nc2;
+ while ((nc2 = OMP_CLAUSE_CHAIN (nc))
+ && OMP_CLAUSE_CODE (nc2) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (nc2) == GOMP_MAP_GRID_STRIDE
+ || OMP_CLAUSE_MAP_KIND (nc2) == GOMP_MAP_SHAPE_DIM))
{
- /* If the element size is not the same as the
- distance between two adjacent array
- elements (in the innermost dimension),
- retrieve the latter value ("span") from the
- size field of the stride. We only expect to
- see one such field per array. */
- gcc_assert (!span);
- span = OMP_CLAUSE_SIZE (nc2);
- span = fold_convert (sizetype, span);
+ if (OMP_CLAUSE_MAP_KIND (nc2) == GOMP_MAP_GRID_STRIDE)
+ {
+ stride = OMP_CLAUSE_DECL (nc2);
+ stride = fold_convert (sizetype, stride);
+ if (OMP_CLAUSE_SIZE (nc2))
+ {
+ /* If the element size is not the same as the
+ distance between two adjacent array
+ elements (in the innermost dimension),
+ retrieve the latter value ("span") from the
+ size field of the stride. We only expect
+ to see one such field per array.
+ ??? Fortran only, e.g. derived-type array ??? */
+ gcc_assert (!span);
+ span = OMP_CLAUSE_SIZE (nc2);
+ span = fold_convert (sizetype, span);
+ }
+ }
+ else if (OMP_CLAUSE_MAP_KIND (nc2) == GOMP_MAP_SHAPE_DIM)
+ dim = OMP_CLAUSE_DECL (nc2);
+ nc = nc2;
}
- nc = nc2;
- }
- if (tdim < tdims.length ())
+ if (dim == NULL_TREE)
+ {
+ if (tdim < tdims.length ()
+ && TREE_CODE (tdims[tdim]) == ARRAY_TYPE)
+ {
+ /* We have an array shape -- use that to find the
+ total size of the data on the target to look up
+ in libgomp:
+ dim = maxval - minval + 1 */
+ tree dtype = TYPE_DOMAIN (tdims[tdim]);
+ tree minval = TYPE_MIN_VALUE (dtype);
+ tree maxval = TYPE_MAX_VALUE (dtype);
+ minval = fold_convert (sizetype, minval);
+ maxval = fold_convert (sizetype, maxval);
+ dim = size_binop (MINUS_EXPR, maxval, minval);
+ dim = size_binop (PLUS_EXPR, dim, size_one_node);
+ }
+ else if ((pointer_based && !handled_pointer_section)
+ || TREE_CODE (tdims[tdim]) == POINTER_TYPE)
+ {
+ /* Use the selected array section to determine the
+ size of the array:
+ dim = index + len * stride - stride + 1 */
+ tree tmp = size_binop (MULT_EXPR, len, stride);
+ tmp = size_binop (MINUS_EXPR, tmp, stride);
+ tmp = size_binop (PLUS_EXPR, tmp, size_one_node);
+ dim = size_binop (PLUS_EXPR, index, tmp);
+ handled_pointer_section = true;
+ }
+ else
+ {
+ if (pointer_based)
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "too many array section specifiers "
+ "for pointer-based array");
+ else
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "too many array section specifiers "
+ "for array");
+ dim = index = len = stride = error_mark_node;
+ }
+ }
+ tdim++;
+
+ c = nc;
+ }
+ else if (TREE_CODE (tdims[tdim]) == ARRAY_TYPE)
{
- /* We have an array shape -- use that to find the
- total size of the data on the target to look up
- in libgomp. */
+ /* We have more array dimensions than array section
+ specifiers. Copy the whole span. */
tree dtype = TYPE_DOMAIN (tdims[tdim]);
tree minval = TYPE_MIN_VALUE (dtype);
tree maxval = TYPE_MAX_VALUE (dtype);
minval = fold_convert (sizetype, minval);
maxval = fold_convert (sizetype, maxval);
dim = size_binop (MINUS_EXPR, maxval, minval);
- dim = size_binop (PLUS_EXPR, dim,
- size_one_node);
- arrsize = size_binop (MULT_EXPR, arrsize, dim);
+ dim = size_binop (PLUS_EXPR, dim, size_one_node);
+ len = dim;
+ index = minval;
+ nc = c;
}
- else if (pointer_based && !handled_pointer_section)
- {
- /* Use the selected array section to determine the
- size of the array. */
- tree tmp = size_binop (MULT_EXPR, len, stride);
- tmp = size_binop (MINUS_EXPR, tmp, stride);
- tmp = size_binop (PLUS_EXPR, tmp, size_one_node);
- dim = size_binop (PLUS_EXPR, index, tmp);
- arrsize = size_binop (MULT_EXPR, arrsize, dim);
- handled_pointer_section = true;
- }
- else
- {
- if (pointer_based)
- error_at (OMP_CLAUSE_LOCATION (c),
- "too many array section specifiers "
- "for pointer-based array");
- else
- error_at (OMP_CLAUSE_LOCATION (c),
- "too many array section specifiers "
- "for array");
- dim = index = len = stride = error_mark_node;
- }
- tdim++;
- c = nc;
+ if (TREE_CODE (dim) != INTEGER_CST)
+ TREE_STATIC (dim_tmp) = 0;
+
+ if (TREE_CODE (index) != INTEGER_CST)
+ TREE_STATIC (index_tmp) = 0;
+
+ if (TREE_CODE (len) != INTEGER_CST)
+ TREE_STATIC (len_tmp) = 0;
+
+ if (TREE_CODE (stride) != INTEGER_CST)
+ TREE_STATIC (stride_tmp) = 0;
+
+ tree cidx = size_int (idim);
+ CONSTRUCTOR_APPEND_ELT (vdim, cidx, dim);
+ CONSTRUCTOR_APPEND_ELT (vindex, cidx, index);
+ CONSTRUCTOR_APPEND_ELT (vlen, cidx, len);
+ CONSTRUCTOR_APPEND_ELT (vstride, cidx, stride);
}
- else
+
+ tree t = size_int (seg_ndims[j]);
+ if (TREE_CODE (t) != INTEGER_CST)
+ TREE_STATIC (seg_ndims_tmp) = 0;
+ CONSTRUCTOR_APPEND_ELT (vseg_ndims, size_int (j), t);
+
+ if (j < segments - 1)
{
- /* We have more array dimensions than array section
- specifiers. Copy the whole span. */
- tree dtype = TYPE_DOMAIN (tdims[tdim]);
- tree minval = TYPE_MIN_VALUE (dtype);
- tree maxval = TYPE_MAX_VALUE (dtype);
- minval = fold_convert (sizetype, minval);
- maxval = fold_convert (sizetype, maxval);
- dim = size_binop (MINUS_EXPR, maxval, minval);
- dim = size_binop (PLUS_EXPR, dim, size_one_node);
- len = dim;
- index = minval;
- nc = c;
+ if (TREE_CODE (nptrs) != INTEGER_CST)
+ TREE_STATIC (seg_nptrs_tmp) = 0;
+ CONSTRUCTOR_APPEND_ELT (vseg_nptrs, size_int (j), nptrs);
}
-
- if (TREE_CODE (dim) != INTEGER_CST)
- TREE_STATIC (dim_tmp) = 0;
-
- if (TREE_CODE (index) != INTEGER_CST)
- TREE_STATIC (index_tmp) = 0;
-
- if (TREE_CODE (len) != INTEGER_CST)
- TREE_STATIC (len_tmp) = 0;
-
- if (TREE_CODE (stride) != INTEGER_CST)
- TREE_STATIC (stride_tmp) = 0;
-
- tree cidx = size_int (i);
- CONSTRUCTOR_APPEND_ELT (vdim, cidx, dim);
- CONSTRUCTOR_APPEND_ELT (vindex, cidx, index);
- CONSTRUCTOR_APPEND_ELT (vlen, cidx, len);
- CONSTRUCTOR_APPEND_ELT (vstride, cidx, stride);
}
tree bias = size_zero_node;
tree volume = size_one_node;
tree enclosure = size_one_node;
- for (i = dims - 1; i >= 0; i--)
+ int last_seg_start = dims - seg_ndims[segments - 1];
+ for (i = dims - 1; i >= last_seg_start; i--)
{
tree dim = (*vdim)[i].value;
tree index = (*vindex)[i].value;
@@ -13942,27 +14059,23 @@ lower_omp_target_grid_desc (tree c, gimple_seq
*ilist_p)
tree index_stride = size_binop (MULT_EXPR, index, stride);
bias = size_binop (PLUS_EXPR, bias,
- size_binop (MULT_EXPR, volume,
- index_stride));
+ size_binop (MULT_EXPR, volume, index_stride));
volume = size_binop (MULT_EXPR, volume, dim);
- if (i == 0)
+ if (i == last_seg_start)
{
- tree elems_covered = size_binop (MINUS_EXPR, len,
- size_one_node);
- elems_covered = size_binop (MULT_EXPR, elems_covered,
- stride);
- elems_covered = size_binop (PLUS_EXPR, elems_covered,
- size_one_node);
- enclosure = size_binop (MULT_EXPR, enclosure,
- elems_covered);
+ /* elems_covered = (len - 1) * stride + 1 */
+ tree elems_covered = size_binop (MINUS_EXPR, len, size_one_node);
+ elems_covered = size_binop (MULT_EXPR, elems_covered, stride);
+ elems_covered = size_binop (PLUS_EXPR, elems_covered, size_one_node);
+ enclosure = size_binop (MULT_EXPR, enclosure, elems_covered);
}
else
enclosure = volume;
}
- /* If we don't have a separate span size, use the element size
- instead. */
+ /* If we don't have a separate span size (??? Fortran only ???), use the
+ element size instead. */
if (!span)
span = fold_convert (sizetype, elsize);
@@ -13970,12 +14083,14 @@ lower_omp_target_grid_desc (tree c, gimple_seq
*ilist_p)
transferred. */
OMP_CLAUSE_SIZE (oc) = size_binop (MULT_EXPR, enclosure, span);
/* And the bias of the first element we will update. */
- OMP_CLAUSE_SIZE (dn) = size_binop (MULT_EXPR, bias, span);
+ OMP_CLAUSE_SIZE (desc_node) = size_binop (MULT_EXPR, bias, span);
- tree cdim = build_constructor (size_arr_type, vdim);
- tree cindex = build_constructor (size_arr_type, vindex);
- tree clen = build_constructor (size_arr_type, vlen);
- tree cstride = build_constructor (size_arr_type, vstride);
+ tree cdim = build_constructor (dim_arr_type, vdim);
+ tree cindex = build_constructor (dim_arr_type, vindex);
+ tree clen = build_constructor (dim_arr_type, vlen);
+ tree cstride = build_constructor (dim_arr_type, vstride);
+ tree cseg_ndims = build_constructor (seg_arr_type, vseg_ndims);
+ tree cseg_nptrs = build_constructor (seg1_arr_type, vseg_nptrs);
if (TREE_STATIC (dim_tmp))
DECL_INITIAL (dim_tmp) = cdim;
@@ -13997,6 +14112,16 @@ lower_omp_target_grid_desc (tree c, gimple_seq
*ilist_p)
else
gimplify_assign (stride_tmp, cstride, ilist_p);
+ if (TREE_STATIC (seg_ndims_tmp))
+ DECL_INITIAL (seg_ndims_tmp) = cseg_ndims;
+ else
+ gimplify_assign (seg_ndims_tmp, cseg_ndims, ilist_p);
+
+ if (TREE_STATIC (seg_nptrs_tmp))
+ DECL_INITIAL (seg_nptrs_tmp) = cseg_nptrs;
+ else
+ gimplify_assign (seg_nptrs_tmp, cseg_nptrs, ilist_p);
+
tree desc_type = TREE_TYPE (desc);
tree ndims_field = TYPE_FIELDS (desc_type);
@@ -14006,35 +14131,47 @@ lower_omp_target_grid_desc (tree c, gimple_seq
*ilist_p)
tree index_field = DECL_CHAIN (dim_field);
tree len_field = DECL_CHAIN (index_field);
tree stride_field = DECL_CHAIN (len_field);
+ tree nsegments_field = DECL_CHAIN (stride_field);
+ tree seg_ndims_field = DECL_CHAIN (nsegments_field);
+ tree seg_nptrs_field = DECL_CHAIN (seg_ndims_field);
vec<constructor_elt, va_gc> *v;
vec_alloc (v, 7);
- bool all_static = (TREE_STATIC (dim_tmp)
- && TREE_STATIC (index_tmp)
- && TREE_STATIC (len_tmp)
- && TREE_STATIC (stride_tmp));
+ bool all_static
+ = (TREE_STATIC (dim_tmp) && TREE_STATIC (index_tmp) && TREE_STATIC
(len_tmp)
+ && TREE_STATIC (stride_tmp) && TREE_STATIC (seg_ndims_tmp)
+ && TREE_STATIC (seg_nptrs_tmp));
- dim_tmp = build4 (ARRAY_REF, sizetype, dim_tmp, size_zero_node,
- NULL_TREE, NULL_TREE);
+ dim_tmp = build4 (ARRAY_REF, sizetype, dim_tmp, size_zero_node, NULL_TREE,
+ NULL_TREE);
dim_tmp = build_fold_addr_expr (dim_tmp);
/* TODO: we could skip all-zeros index. */
- index_tmp = build4 (ARRAY_REF, sizetype, index_tmp,
- size_zero_node, NULL_TREE, NULL_TREE);
+ index_tmp = build4 (ARRAY_REF, sizetype, index_tmp, size_zero_node,
NULL_TREE,
+ NULL_TREE);
index_tmp = build_fold_addr_expr (index_tmp);
- len_tmp = build4 (ARRAY_REF, sizetype, len_tmp, size_zero_node,
- NULL_TREE, NULL_TREE);
+ len_tmp = build4 (ARRAY_REF, sizetype, len_tmp, size_zero_node, NULL_TREE,
+ NULL_TREE);
len_tmp = build_fold_addr_expr (len_tmp);
/* TODO: we could skip all-ones stride. */
- stride_tmp = build4 (ARRAY_REF, sizetype, stride_tmp,
- size_zero_node, NULL_TREE, NULL_TREE);
+ stride_tmp = build4 (ARRAY_REF, sizetype, stride_tmp, size_zero_node,
+ NULL_TREE, NULL_TREE);
stride_tmp = build_fold_addr_expr (stride_tmp);
+ seg_ndims_tmp = build4 (ARRAY_REF, sizetype, seg_ndims_tmp, size_zero_node,
+ NULL_TREE, NULL_TREE);
+ seg_ndims_tmp = build_fold_addr_expr (seg_ndims_tmp);
+
+ seg_nptrs_tmp = build4 (ARRAY_REF, sizetype, seg_nptrs_tmp, size_zero_node,
+ NULL_TREE, NULL_TREE);
+ seg_nptrs_tmp = build_fold_addr_expr (seg_nptrs_tmp);
+
elsize = fold_convert (sizetype, elsize);
tree ndims = size_int (dims);
+ tree nsegments = size_int (segments);
CONSTRUCTOR_APPEND_ELT (v, ndims_field, ndims);
CONSTRUCTOR_APPEND_ELT (v, elemsize_field, elsize);
@@ -14043,6 +14180,9 @@ lower_omp_target_grid_desc (tree c, gimple_seq *ilist_p)
CONSTRUCTOR_APPEND_ELT (v, index_field, index_tmp);
CONSTRUCTOR_APPEND_ELT (v, len_field, len_tmp);
CONSTRUCTOR_APPEND_ELT (v, stride_field, stride_tmp);
+ CONSTRUCTOR_APPEND_ELT (v, nsegments_field, nsegments);
+ CONSTRUCTOR_APPEND_ELT (v, seg_ndims_field, seg_ndims_tmp);
+ CONSTRUCTOR_APPEND_ELT (v, seg_nptrs_field, seg_nptrs_tmp);
tree desc_ctor = build_constructor (desc_type, v);
@@ -14054,9 +14194,7 @@ lower_omp_target_grid_desc (tree c, gimple_seq *ilist_p)
else
gimplify_assign (desc, desc_ctor, ilist_p);
- OMP_CLAUSE_CHAIN (dn) = OMP_CLAUSE_CHAIN (nc);
-}
-
+ OMP_CLAUSE_CHAIN (desc_node) = OMP_CLAUSE_CHAIN (nc);
}
/* Lower the GIMPLE_OMP_TARGET in the current statement
@@ -14208,6 +14346,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
case GOMP_MAP_FROM_GRID:
case GOMP_MAP_GRID_DIM:
case GOMP_MAP_GRID_STRIDE:
+ case GOMP_MAP_SHAPE_DIM:
break;
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_FORCE_ALLOC:
@@ -14245,7 +14384,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
while ((nc = OMP_CLAUSE_CHAIN (c))
&& OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM
- || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE))
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_SHAPE_DIM))
c = nc;
map_cnt += 2;
continue;
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-1.c
b/gcc/testsuite/c-c++-common/gomp/array-section-1.c
index 99756bbf4c2..e787466d1d1 100644
--- a/gcc/testsuite/c-c++-common/gomp/array-section-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-1.c
@@ -13,6 +13,13 @@ 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" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_dim\.[0-9]+\[2\] = [{]5, 10[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_index\.[0-9]+\[2\] = [{]2, 0[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_len\.[0-9]+\[2\] = [{]1, 10[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_stride\.[0-9]+\[2\] = [{]1, 1[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_seg_ndims\.[0-9]+\[2\] = [{]1, 1[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_seg_nptrs\.[0-9]+\[1\] = [{]1[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static struct \.omp_nc_desc_x\.[0-9]+ =
[{]\.__ndims=2, \.__elemsize=4, \.__span=4, \.__dim=&\.omp_dim\.[0-9]+\[0\],
\.__index=&\.omp_index\.[0-9]+\[0\], \.__length=&\.omp_len\.[0-9]+\[0\],
\.__stride=&\.omp_stride\.[0-9]+\[0\], \.__nsegments=2,
\.__seg_ndims=&\.omp_seg_ndims\.[0-9]+\[0\],
\.__seg_nptrs=&\.omp_seg_nptrs\.[0-9]+\[0\][}];} "lower" } } */
}
/* Index, length and stride need not be literal constants -- a variable
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-2.c
b/gcc/testsuite/c-c++-common/gomp/array-section-2.c
index 15bc2cb28ba..271b5ce31b5 100644
--- a/gcc/testsuite/c-c++-common/gomp/array-section-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-2.c
@@ -12,4 +12,11 @@ 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" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_dim\.[0-9]+\[2\] = [{]5, 5[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_index\.[0-9]+\[2\] = [{]0, 0[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_len\.[0-9]+\[2\] = [{]1, 2[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_stride\.[0-9]+\[2\] = [{]1, 1[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_seg_ndims\.[0-9]+\[1\] = [{]2[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_seg_nptrs\.[0-9]+\[1\] = [{][}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static struct \.omp_nc_desc_y\.[0-9]+ =
[{]\.__ndims=2, \.__elemsize=8, \.__span=8, \.__dim=&\.omp_dim\.[0-9]+\[0\],
\.__index=&\.omp_index\.[0-9]+\[0\], \.__length=&\.omp_len\.[0-9]+\[0\],
\.__stride=&\.omp_stride\.[0-9]+\[0\], \.__nsegments=1,
\.__seg_ndims=&\.omp_seg_ndims\.[0-9]+\[0\],
\.__seg_nptrs=&\.omp_seg_nptrs\.[0-9]+\[0\][}];} "lower" } } */
}
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-3.c
b/gcc/testsuite/c-c++-common/gomp/array-section-3.c
index f7b17475b5d..f4a35e839db 100644
--- a/gcc/testsuite/c-c++-common/gomp/array-section-3.c
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-3.c
@@ -16,6 +16,13 @@ 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" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_dim\.[0-9]+\[4\] = [{]3, 4, 6, 6[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_index\.[0-9]+\[4\] = [{]1, 1, 1, 0[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_len\.[0-9]+\[4\] = [{]1, 1, 3, 2[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_stride\.[0-9]+\[4\] = [{]1, 1, 2, 1[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_seg_ndims\.[0-9]+\[2\] = [{]2, 2[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_seg_nptrs\.[0-9]+\[1\] = [{]1[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static struct \.omp_nc_desc_x\.[0-9]+ =
[{]\.__ndims=4, \.__elemsize=4, \.__span=4, \.__dim=&\.omp_dim\.[0-9]+\[0\],
\.__index=&\.omp_index\.[0-9]+\[0\], \.__length=&\.omp_len\.[0-9]+\[0\],
\.__stride=&\.omp_stride\.[0-9]+\[0\], \.__nsegments=2,
\.__seg_ndims=&\.omp_seg_ndims\.[0-9]+\[0\],
\.__seg_nptrs=&\.omp_seg_nptrs\.[0-9]+\[0\][}];} "lower" } } */
}
void shape_cast_length_one_range_index (void)
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-4.c
b/gcc/testsuite/c-c++-common/gomp/array-section-4.c
index ce446b931c2..f9a9e8866aa 100644
--- a/gcc/testsuite/c-c++-common/gomp/array-section-4.c
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-4.c
@@ -18,4 +18,11 @@ 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" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_dim\.[0-9]+\[4\] = [{]3, 3, 4, 6[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_index\.[0-9]+\[4\] = [{]1, 2, 0, 2[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_len\.[0-9]+\[4\] = [{]1, 1, 1, 4[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_stride\.[0-9]+\[4\] = [{]1, 1, 1, 1[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_seg_ndims\.[0-9]+\[2\] = [{]3, 1[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_seg_nptrs\.[0-9]+\[1\] = [{]1[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static struct \.omp_nc_desc_w\.[0-9]+ =
[{]\.__ndims=4, \.__elemsize=4, \.__span=4, \.__dim=&\.omp_dim\.[0-9]+\[0\],
\.__index=&\.omp_index\.[0-9]+\[0\], \.__length=&\.omp_len\.[0-9]+\[0\],
\.__stride=&\.omp_stride\.[0-9]+\[0\], \.__nsegments=2,
\.__seg_ndims=&\.omp_seg_ndims\.[0-9]+\[0\],
\.__seg_nptrs=&\.omp_seg_nptrs\.[0-9]+\[0\][}];} "lower" } } */
}
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-5.c
b/gcc/testsuite/c-c++-common/gomp/array-section-5.c
index 9bd8a251393..bd136ead018 100644
--- a/gcc/testsuite/c-c++-common/gomp/array-section-5.c
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-5.c
@@ -15,6 +15,13 @@ 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" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_dim\.[0-9]+\[4\] = [{]3, 4, 6, 6[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_index\.[0-9]+\[4\] = [{]1, 1, 1, 0[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_len\.[0-9]+\[4\] = [{]1, 1, 3, 6[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_stride\.[0-9]+\[4\] = [{]1, 1, 2, 1[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_seg_ndims\.[0-9]+\[2\] = [{]2, 2[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_seg_nptrs\.[0-9]+\[1\] = [{]1[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static struct \.omp_nc_desc_x\.[0-9]+ =
[{]\.__ndims=4, \.__elemsize=4, \.__span=4, \.__dim=&\.omp_dim\.[0-9]+\[0\],
\.__index=&\.omp_index\.[0-9]+\[0\], \.__length=&\.omp_len\.[0-9]+\[0\],
\.__stride=&\.omp_stride\.[0-9]+\[0\], \.__nsegments=2,
\.__seg_ndims=&\.omp_seg_ndims\.[0-9]+\[0\],
\.__seg_nptrs=&\.omp_seg_nptrs\.[0-9]+\[0\][}];} "lower" } } */
}
void shape_cast_omitted_column_dim_length_one_range (void)
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-6.c
b/gcc/testsuite/c-c++-common/gomp/array-section-6.c
index e90f4e66188..0199f725eda 100644
--- a/gcc/testsuite/c-c++-common/gomp/array-section-6.c
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-6.c
@@ -15,4 +15,5 @@ void plain_pointer_shape_cast (void)
#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" } } */
+/* { dg-final { scan-tree-dump-not "__ndims" "lower" } } */
}
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-8.c
b/gcc/testsuite/c-c++-common/gomp/array-section-8.c
index 4873138362a..986d774782f 100644
--- a/gcc/testsuite/c-c++-common/gomp/array-section-8.c
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-8.c
@@ -12,6 +12,13 @@ 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" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_dim\.[0-9]+\[3\] = [{]4, 3, 12[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_index\.[0-9]+\[3\] = [{]2, 0, 3[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_len\.[0-9]+\[3\] = [{]1, 2, 5[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_stride\.[0-9]+\[3\] = [{]1, 2, 2[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_seg_ndims\.[0-9]+\[3\] = [{]1, 1, 1[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_seg_nptrs\.[0-9]+\[2\] = [{]1, 2[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static struct \.omp_nc_desc_arr\.[0-9]+ =
[{]\.__ndims=3, \.__elemsize=4, \.__span=4, \.__dim=&\.omp_dim\.[0-9]+\[0\],
\.__index=&\.omp_index\.[0-9]+\[0\], \.__length=&\.omp_len\.[0-9]+\[0\],
\.__stride=&\.omp_stride\.[0-9]+\[0\], \.__nsegments=3,
\.__seg_ndims=&\.omp_seg_ndims\.[0-9]+\[0\],
\.__seg_nptrs=&\.omp_seg_nptrs\.[0-9]+\[0\][}];} "lower" } } */
}
/* Same shape, but every index/length/stride across all three segments is
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-9.c
b/gcc/testsuite/c-c++-common/gomp/array-section-9.c
index 9d2e0ceae09..eeecb1faa57 100644
--- a/gcc/testsuite/c-c++-common/gomp/array-section-9.c
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-9.c
@@ -18,4 +18,11 @@ 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" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_dim\.[0-9]+\[6\] = [{]2, 3, 3, 3, 2, 20[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_index\.[0-9]+\[6\] = [{]1, 0, 0, 0, 0, 3[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_len\.[0-9]+\[6\] = [{]1, 2, 2, 2, 2, 5[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_stride\.[0-9]+\[6\] = [{]1, 2, 2, 2, 1, 2[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_seg_ndims\.[0-9]+\[3\] = [{]2, 2, 2[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static long unsigned int
\.omp_seg_nptrs\.[0-9]+\[2\] = [{]2, 4[}];} "lower" } } */
+/* { dg-final { scan-tree-dump {static struct \.omp_nc_desc_arr\.[0-9]+ =
[{]\.__ndims=6, \.__elemsize=4, \.__span=4, \.__dim=&\.omp_dim\.[0-9]+\[0\],
\.__index=&\.omp_index\.[0-9]+\[0\], \.__length=&\.omp_len\.[0-9]+\[0\],
\.__stride=&\.omp_stride\.[0-9]+\[0\], \.__nsegments=3,
\.__seg_ndims=&\.omp_seg_ndims\.[0-9]+\[0\],
\.__seg_nptrs=&\.omp_seg_nptrs\.[0-9]+\[0\][}];} "lower" } } */
}
--
2.53.0
