From: Julian Brown <[email protected]>
This patch implements noncontiguous "target update" for Fortran.
The existing middle end/runtime bits relating to C++ support are reused,
with some small adjustments, e.g.:
1. The node used to map the OMP "array descriptor" (from omp-low.cc
onwards) now uses the OMP_CLAUSE_SIZE field as a bias (the difference
between the "virtual origin" element with zero indices in each
dimension and the first element actually stored in memory).
2. The OMP_CLAUSE_SIZE field of a GOMP_MAP_DIM_STRIDE node may now be
used to store a "span", which is the distance in bytes between
two adjacent elements in an array (with unit stride) when that is
different from the element size, as it can be in Fortran.
The implementation goes to some effort to massage Fortran array metadata
(array descriptors) into a form that can ultimately be consumed by
omp_target_memcpy_rect_worker. The method for doing this is described
in comments in the patch body.
As with C and C++, strided updates in combination with map iterators are
not (yet) supported. However, contrary to these languages, Fortran does not
permit to detect non-unit strides statically. Therefore, a runtime check
had to be added.
The middle-end/runtime "bias" and "span" support this patch relies on
(and the VIEW_CONVERT_EXPR/OMP_CLAUSE_MAP fixes it needed along the way)
is not actually Fortran-specific, so it has been folded into the earlier
C++ patch instead; only the Fortran front end and its tests remain here.
2023-07-03 Julian Brown <[email protected]>
gcc/fortran/ChangeLog
* trans-openmp.cc (gfc_omp_deep_map_kind_p): Handle
GOMP_MAP_{TO,FROM}_GRID, GOMP_MAP_GRID_{DIM,STRIDE}.
(gfc_trans_omp_arrayshape_type, gfc_omp_calculate_gcd,
gfc_desc_to_omp_noncontig_array, gfc_omp_contiguous_update_p): New
functions.
(gfc_trans_omp_clauses): Handle noncontiguous to/from clauses for OMP
"target update" directives.
gcc/testsuite/ChangeLog
* gfortran.dg/gomp/noncontig-updates-1.f90: New test.
* gfortran.dg/gomp/noncontig-updates-2.f90: New test.
* gfortran.dg/gomp/noncontig-updates-3.f90: New test.
* gfortran.dg/gomp/noncontig-updates-4.f90: New test.
libgomp/ChangeLog
* testsuite/libgomp.fortran/noncontig-updates-1.f90: New test.
* testsuite/libgomp.fortran/noncontig-updates-2.f90: New test.
* testsuite/libgomp.fortran/noncontig-updates-3.f90: New test.
* testsuite/libgomp.fortran/noncontig-updates-4.f90: New test.
* testsuite/libgomp.fortran/noncontig-updates-5.f90: New test.
* testsuite/libgomp.fortran/noncontig-updates-6.f90: New test.
* testsuite/libgomp.fortran/noncontig-updates-7.f90: New test.
* testsuite/libgomp.fortran/noncontig-updates-8.f90: New test.
* testsuite/libgomp.fortran/noncontig-updates-9.f90: New test.
* testsuite/libgomp.fortran/noncontig-updates-10.f90: New test.
* testsuite/libgomp.fortran/noncontig-updates-11.f90: New test.
* testsuite/libgomp.fortran/noncontig-updates-12.f90: New test.
* testsuite/libgomp.fortran/noncontig-updates-13.f90: New test.
2026-08-13 Paul-Antoine Arras <[email protected]>
gcc/fortran/
* openmp.cc (resolve_omp_clauses_aff_dep_map_cache): Exempt
EXEC_OMP_TARGET_UPDATE from the contiguity/no-stride checks,
like EXEC_OACC_UPDATE.
* trans-openmp.cc (gfc_trans_omp_clauses): Route the
noncontiguous-grid case's statements through the iterator-scoped
&iter_block instead of the outer block. Take the regular
map-clause path when an iterator is present, with a runtime
contiguity check.
* expr.cc (gfc_is_simply_contiguous): Also return true for a
whole ref chain of only element-selecting array refs.
gcc/testsuite/
* gfortran.dg/gomp/target-update-iterators-3.f90: Update
scan-dump patterns.
* gfortran.dg/gomp/target-update-iterators-4.f90: New test.
libgomp/testsuite/
* libgomp.fortran/target-update-iterators-4.f90: New test.
2026-09-25 Paul-Antoine Arras <[email protected]>
gcc/fortran/
* openmp.cc (resolve_omp_clauses_aff_dep_map_cache): Reject a
map iterator combined with a non-contiguous array or a non-unit
stride. Restrict the constant positive-stride check to a
genuinely constant integer stride. Remove a stray trailing
space.
gcc/testsuite/
* gfortran.dg/gomp/bad-array-section-1.f90: New test.
* gfortran.dg/gomp/noncontig-updates-5.f90: New test.
* gfortran.dg/gomp/noncontig-updates-6.f90: New test.
* gfortran.dg/gomp/target-update-iterators-3.f90: Fix expected
diagnostic text to match the hyphenated "non-contiguous"
wording.
* gfortran.dg/gomp/target-update-iterators-5.f90: New test.
* gfortran.dg/gomp/target-update-iterators-6.f90: New test.
libgomp/testsuite/
* libgomp.fortran/target-update-iterators-5.f90: New test.
2026-10-16 Paul-Antoine Arras <[email protected]>
gcc/fortran/
* openmp.cc (resolve_omp_clauses_aff_dep_map_cache): Only stop at
the first strided dimension on error, so that later dimensions are
checked too.
* trans-openmp.cc (gfc_trans_omp_clauses): Use the rank of the
expression as the number of dimensions of a noncontiguous update,
and also handle strided sections of components. Take the span from
the descriptor of a pointer to a component.
gcc/testsuite/
* gfortran.dg/gomp/noncontig-updates-7.f90: New test.
* gfortran.dg/gomp/noncontig-updates-8.f90: New test.
* gfortran.dg/gomp/noncontig-updates-9.f90: New test.
libgomp/testsuite/
* libgomp.fortran/noncontig-updates-14.f90: New test.
* libgomp.fortran/noncontig-updates-15.f90: New test.
* libgomp.fortran/noncontig-updates-16.f90: New test.
* libgomp.fortran/noncontig-updates-17.f90: New test.
* libgomp.fortran/noncontig-updates-18.f90: New test.
* libgomp.fortran/target-update-iterators-5.f90: XFAIL for PR 127616.
Co-Authored-By: Sandra Loosemore <[email protected]>
---
gcc/fortran/expr.cc | 3 +
gcc/fortran/openmp.cc | 29 +-
gcc/fortran/trans-openmp.cc | 522 ++++++++++++++++++
.../gfortran.dg/gomp/bad-array-section-1.f90 | 31 ++
.../gfortran.dg/gomp/noncontig-updates-1.f90 | 19 +
.../gfortran.dg/gomp/noncontig-updates-2.f90 | 16 +
.../gfortran.dg/gomp/noncontig-updates-3.f90 | 16 +
.../gfortran.dg/gomp/noncontig-updates-4.f90 | 15 +
.../gfortran.dg/gomp/noncontig-updates-5.f90 | 28 +
.../gfortran.dg/gomp/noncontig-updates-6.f90 | 15 +
.../gfortran.dg/gomp/noncontig-updates-7.f90 | 17 +
.../gfortran.dg/gomp/noncontig-updates-8.f90 | 37 ++
.../gfortran.dg/gomp/noncontig-updates-9.f90 | 31 ++
.../gomp/target-update-iterators-3.f90 | 11 +-
.../gomp/target-update-iterators-4.f90 | 21 +
.../gomp/target-update-iterators-5.f90 | 23 +
.../gomp/target-update-iterators-6.f90 | 39 ++
.../libgomp.fortran/noncontig-updates-1.f90 | 54 ++
.../libgomp.fortran/noncontig-updates-10.f90 | 29 +
.../libgomp.fortran/noncontig-updates-11.f90 | 51 ++
.../libgomp.fortran/noncontig-updates-12.f90 | 59 ++
.../libgomp.fortran/noncontig-updates-13.f90 | 42 ++
.../libgomp.fortran/noncontig-updates-14.f90 | 41 ++
.../libgomp.fortran/noncontig-updates-15.f90 | 95 ++++
.../libgomp.fortran/noncontig-updates-16.f90 | 163 ++++++
.../libgomp.fortran/noncontig-updates-17.f90 | 36 ++
.../libgomp.fortran/noncontig-updates-18.f90 | 150 +++++
.../libgomp.fortran/noncontig-updates-2.f90 | 101 ++++
.../libgomp.fortran/noncontig-updates-3.f90 | 47 ++
.../libgomp.fortran/noncontig-updates-4.f90 | 78 +++
.../libgomp.fortran/noncontig-updates-5.f90 | 55 ++
.../libgomp.fortran/noncontig-updates-6.f90 | 34 ++
.../libgomp.fortran/noncontig-updates-7.f90 | 36 ++
.../libgomp.fortran/noncontig-updates-8.f90 | 39 ++
.../libgomp.fortran/noncontig-updates-9.f90 | 34 ++
.../target-update-iterators-4.f90 | 36 ++
.../target-update-iterators-5.f90 | 56 ++
37 files changed, 2106 insertions(+), 3 deletions(-)
create mode 100644 gcc/testsuite/gfortran.dg/gomp/bad-array-section-1.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/noncontig-updates-1.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/noncontig-updates-2.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/noncontig-updates-3.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/noncontig-updates-4.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/noncontig-updates-5.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/noncontig-updates-6.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/noncontig-updates-7.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/noncontig-updates-8.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/noncontig-updates-9.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-update-iterators-4.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-update-iterators-5.f90
create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-update-iterators-6.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-1.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-10.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-11.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-12.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-13.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-14.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-15.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-16.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-17.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-18.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-2.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-3.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-4.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-5.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-6.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-7.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-8.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/noncontig-updates-9.f90
create mode 100644
libgomp/testsuite/libgomp.fortran/target-update-iterators-4.f90
create mode 100644
libgomp/testsuite/libgomp.fortran/target-update-iterators-5.f90
diff --git a/gcc/fortran/expr.cc b/gcc/fortran/expr.cc
index 52945a1db13..362d083c5bc 100644
--- a/gcc/fortran/expr.cc
+++ b/gcc/fortran/expr.cc
@@ -6474,6 +6474,9 @@ gfc_is_simply_contiguous (gfc_expr *expr, bool strict,
bool permit_element)
ar = &ref->u.ar;
}
+ if (permit_element && !ar && !part_ref)
+ return true;
+
sym = expr->symtree->n.sym;
if ((part_ref
&& part_ref->u.c.component
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index a2c825401e3..9b789e74a6d 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -10409,6 +10409,7 @@ resolve_omp_clauses_aff_dep_map_cache (gfc_code *code,
could be contiguous even if it looks like it may not be. */
if (code
&& code->op != EXEC_OACC_UPDATE
+ && code->op != EXEC_OMP_TARGET_UPDATE
&& list != OMP_LIST_CACHE
&& list != OMP_LIST_DEPEND
&& !gfc_is_simply_contiguous (n->expr, false, true)
@@ -10416,6 +10417,12 @@ resolve_omp_clauses_aff_dep_map_cache (gfc_code *code,
&& !(lastslice && (lastslice->next
|| lastslice->type != REF_ARRAY)))
gfc_error ("Array is not contiguous at %L", &n->where);
+ else if (code && code->op == EXEC_OMP_TARGET_UPDATE
+ && gfc_is_not_contiguous (n->expr) && n->u2.ns)
+ {
+ gfc_error ("Non-contiguous array with map iterator at %L",
+ &n->where);
+ }
}
}
if (list == OMP_LIST_MAP
@@ -10460,12 +10467,32 @@ resolve_omp_clauses_aff_dep_map_cache (gfc_code *code,
int i;
gfc_array_ref *ar = &lastslice->u.ar;
for (i = 0; i < ar->dimen; i++)
- if (ar->stride[i] && code && code->op != EXEC_OACC_UPDATE)
+ if (ar->stride[i] && code && code->op != EXEC_OACC_UPDATE
+ && code->op != EXEC_OMP_TARGET_UPDATE)
{
gfc_error ("Stride should not be specified for array section "
"in %s clause at %L", name, &n->where);
break;
}
+ else if (ar->stride[i] && ar->stride[i]->expr_type == EXPR_CONSTANT
+ && ar->stride[i]->ts.type == BT_INTEGER && code
+ && code->op == EXEC_OMP_TARGET_UPDATE)
+ {
+ if (mpz_sgn (ar->stride[i]->value.integer) <= 0)
+ {
+ gfc_error ("Stride must be positive in %s clause at %L",
+ name, &ar->stride[i]->where);
+ break;
+ }
+ else if (mpz_cmp_si (ar->stride[i]->value.integer, 1) != 0
+ && n->u2.ns)
+ {
+ gfc_error (
+ "Non-unit stride with iterator in %s clause at %L", name,
+ &ar->stride[i]->where);
+ break;
+ }
+ }
else if (ar->dimen_type[i] != DIMEN_ELEMENT
&& ar->dimen_type[i] != DIMEN_RANGE)
{
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 3e4375f6d60..874b54f377d 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -2433,6 +2433,10 @@ gfc_omp_deep_map_kind_p (tree clause)
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ATTACH_DETACH:
+ case GOMP_MAP_TO_GRID:
+ case GOMP_MAP_FROM_GRID:
+ case GOMP_MAP_GRID_DIM:
+ case GOMP_MAP_GRID_STRIDE:
break;
default:
gcc_unreachable ();
@@ -3919,6 +3923,346 @@ gfc_map_array_descriptor (
return false;
}
+/* We build an "un-Fortrannish" array-of-arrays here to pass our calculated
+ array bounds to the middle end for strided/rectangular OpenMP
+ "target update" operations. */
+
+static tree
+gfc_trans_omp_arrayshape_type (tree type, vec<tree> *dims)
+{
+ gcc_assert (dims->length () > 0);
+
+ for (unsigned i = 0; i < dims->length (); i++)
+ {
+ tree dim = fold_convert (sizetype, (*dims)[i]);
+ /* We need the index of the last element, not the array size. */
+ dim = size_binop (MINUS_EXPR, dim, size_one_node);
+ tree idxtype = build_index_type (dim);
+ type = build_array_type (type, idxtype);
+ }
+
+ return type;
+}
+
+/* Emit code to find the greatest common divisor of two (gfc_array_index_type)
+ trees to BLOCK. This is Euclid's algorithm:
+
+ int
+ gcd (int a, int b)
+ {
+ int tmp;
+ while (b != 0)
+ {
+ tmp = b;
+ b = a % b;
+ a = tmp;
+ }
+ return a;
+ }
+*/
+
+static void
+gfc_omp_calculate_gcd (stmtblock_t *block, tree dst, tree a, tree b)
+{
+ tree tmp = gfc_create_var (gfc_array_index_type, "tmp");
+ tree avar = gfc_create_var (gfc_array_index_type, "a");
+ tree bvar = gfc_create_var (gfc_array_index_type, "b");
+
+ /* Avoid clobbering the inputs. */
+ gfc_add_modify (block, avar, a);
+ gfc_add_modify (block, bvar, b);
+
+ tree label_cond = gfc_build_label_decl (NULL_TREE);
+ tree label_loop = gfc_build_label_decl (NULL_TREE);
+ TREE_USED (label_cond) = 1;
+ TREE_USED (label_loop) = 1;
+
+ gfc_add_expr_to_block (block, build1_v (GOTO_EXPR, label_cond));
+ gfc_add_expr_to_block (block, build1_v (LABEL_EXPR, label_loop));
+
+ gfc_add_modify (block, tmp, bvar);
+ gfc_add_modify (block, bvar,
+ fold_build2_loc (input_location, TRUNC_MOD_EXPR,
+ gfc_array_index_type, avar, bvar));
+ gfc_add_modify (block, avar, tmp);
+
+ gfc_add_expr_to_block (block, build1_v (LABEL_EXPR, label_cond));
+
+ tmp = fold_build2_loc (input_location, NE_EXPR, boolean_type_node, bvar,
+ gfc_index_zero_node);
+ tmp = build3_v (COND_EXPR, tmp, build1_v (GOTO_EXPR, label_loop),
+ build_empty_stmt (input_location));
+ gfc_add_expr_to_block (block, tmp);
+
+ gfc_add_modify (block, dst, avar);
+}
+
+/* Convert a gfortran array descriptor -- specifically the per-dimension
+ strides -- into a form that can be easily translated to a noncontiguous
+ OpenMP "target update" operation. We emit a specialized version of a
+ function like this inline:
+
+ void
+ gfc_desc_to_omp_noncontig_array (int *dims, int *strides, int ndims,
+ int *fstrides, int *flo, int *fhi)
+ {
+ dims[ndims - 1] = (fhi[ndims - 1] - flo[ndims - 1] + 1);
+ strides[0] = fstrides[0];
+ if (ndims > 1)
+ strides[ndims - 1] = 1;
+ if (ndims == 2)
+ dims[0] = fstrides[1];
+ else if (ndims > 2)
+ {
+ int grains[ndims - 2];
+
+ int bigger_grain = fstrides[ndims - 1];
+ for (int i = ndims - 2; i > 0; i--)
+ {
+ grains[i - 1] = gcd (fstrides[i], bigger_grain);
+ bigger_grain = grains[i - 1];
+ }
+
+ int volume = 1;
+ for (int i = 0; i < ndims - 2; i++)
+ {
+ int g = grains[i];
+ dims[i] = g / volume;
+ strides[i + 1] = fstrides[i + 1] / g;
+ volume = volume * dims[i];
+ }
+ dims[ndims - 2] = fstrides[ndims - 1] / volume;
+ }
+ }
+
+ where "fstrides", "flo" and "fhi" represent the stride, low bound and upper
+ bound of each dimension in the Fortran array descriptor.
+
+ (Note that most of the complexity only applies to arrays with more than two
+ dimensions, and the final stanza won't be emitted at all for lower-ranked
+ arrays.)
+
+ The output of the algorithm is a set of dimensions dims[] = { D, C, B, A }
+ "as if" the array was declared like this (in C):
+
+ type arr[A][B][C][D];
+
+ i.e. with the innermost dimension first, and a set of strides (in terms of
+ the step size along each dimension, without previous dimensions multiplied
+ in).
+
+ As an example, if we have an array:
+
+ allocate (arr(18,19,20,21,22))
+
+ and an update operation:
+
+ !$omp target update to(arr(1:3:2,1:4:3,1:5:4,1:6:5,1:7:6))
+
+ the strides we see in the Fortran array descriptor will be:
+
+ 2 54 1368 34200 861840
+
+ as given by:
+
+ 2 = stride0
+ 54 = dim0 * stride1
+ 1368 = dim0 * dim1 * stride2
+ 34200 = dim0 * dim1 * dim2 * stride3
+ 861840 = dim0 * dim1 * dim2 * dim3 * stride4
+
+ where "dimN" are the extents of each dimension (18,19,20,21,22), and
+ "strideN" are the strides in terms of step length along each dimension
+ (2,3,4,5,6).
+
+ We'd like to figure out what the original dimN, strideN were from the
+ Fortran array descriptor, but that's in general impossible. Furthermore,
+ if we naively divide a stride by the preceding stride, the result isn't
+ necessarily an integer, as for e.g.:
+
+ 861840/34200 = 25.2
+
+ What we can do though is figure out the greatest common divisor of
+ each stride and the preceding one, from the largest down, and use those as
+ units of granularity, i.e. the size of the corresponding dimension we pass
+ to the middle-end/runtime. The stepwise stride is then the number of
+ times each "grain" fits into the Fortran array descriptor stride.
+
+ The output of the algorithm will be:
+
+ dims strides
+ 18 2
+ 76 3
+ 5 1
+ 126 5
+ 9 1
+
+ These numbers work fine for libgomp target.c:omp_target_memcpy_rect_worker.
+ Multiplying them through also gives the same numbers as the source Fortran
+ array strides, i.e. dim0*dim1*dim2*stride3 (18*76*5*5) = 34200. */
+
+static void
+gfc_desc_to_omp_noncontig_array (stmtblock_t *block, tree *ompdimsp,
+ tree *ompstridesp, tree desc, int ndims)
+{
+ tree lastdim = build_int_cst (gfc_array_index_type, ndims - 1);
+ tree dimrange = build_index_type (lastdim);
+ tree ndimarrtype = build_array_type (gfc_array_index_type, dimrange);
+ tree ompdims = gfc_create_var (ndimarrtype, "dims");
+ tree ompstrides = gfc_create_var (ndimarrtype, "strides");
+
+ *ompdimsp = ompdims;
+ *ompstridesp = ompstrides;
+
+ /* dims[ndims - 1] = (fhi[ndims - 1] - flo[ndims - 1] + 1); */
+ tree lastlbound = gfc_conv_array_lbound (desc, ndims - 1);
+ tree lastubound = gfc_conv_array_ubound (desc, ndims - 1);
+ tree lastrange = fold_build2_loc (input_location, MINUS_EXPR,
+ gfc_array_index_type, lastubound,
+ lastlbound);
+ lastrange = fold_build2_loc (input_location, PLUS_EXPR, gfc_array_index_type,
+ lastrange, gfc_index_one_node);
+
+ gfc_add_modify (block,
+ gfc_build_array_ref (ompdims, lastdim, NULL_TREE, true),
+ lastrange);
+
+ /* strides[0] = fstrides[0]; */
+ tree stride0 = gfc_conv_array_stride (desc, 0);
+ gfc_add_modify (block,
+ gfc_build_array_ref (ompstrides, gfc_index_zero_node,
+ NULL_TREE, true),
+ stride0);
+
+ if (ndims > 1)
+ /* strides[ndims - 1] = 1; */
+ gfc_add_modify (block,
+ gfc_build_array_ref (ompstrides, lastdim, NULL_TREE, true),
+ gfc_index_one_node);
+
+ if (ndims == 2)
+ /* dims[0] = fstrides[1]; */
+ gfc_add_modify (block,
+ gfc_build_array_ref (ompdims, gfc_index_zero_node,
+ NULL_TREE, true),
+ gfc_conv_array_stride (desc, 1));
+ else if (ndims > 2)
+ {
+ /* int grains[ndims - 2]; */
+ tree lastgrain = build_int_cst (gfc_array_index_type, ndims - 3);
+ tree grainrange = build_index_type (lastgrain);
+ tree grainarrtype = build_array_type (gfc_array_index_type, grainrange);
+ tree grains = gfc_create_var (grainarrtype, "grains");
+
+ /* int bigger_grain = fstrides[ndims - 1]; */
+ tree bigger_grain = gfc_create_var (gfc_array_index_type,
"bigger_grain");
+ tree fstridem1 = gfc_conv_array_stride (desc, ndims - 1);
+ gfc_add_modify (block, bigger_grain, fstridem1);
+
+ /*
+ for (int i = ndims - 2; i > 0; i--)
+ {
+ grains[i - 1] = gcd (fstrides[i], bigger_grain);
+ bigger_grain = grains[i - 1];
+ }
+ */
+ stmtblock_t loop_body;
+ gfc_init_block (&loop_body);
+
+ tree idx = gfc_create_var (gfc_array_index_type, "idx");
+
+ tree gcdtmp = gfc_create_var (gfc_array_index_type, "tmp");
+ gfc_omp_calculate_gcd (&loop_body, gcdtmp,
+ gfc_conv_descriptor_stride_get (desc, idx),
+ bigger_grain);
+ tree idxm1 = fold_build2_loc (input_location, MINUS_EXPR,
+ gfc_array_index_type, idx,
+ gfc_index_one_node);
+ gfc_add_modify (&loop_body,
+ gfc_build_array_ref (grains, idxm1, NULL_TREE, true),
+ gcdtmp);
+ gfc_add_modify (&loop_body, bigger_grain, gcdtmp);
+
+ gfc_simple_for_loop (block, idx,
+ build_int_cst (gfc_array_index_type, ndims - 2),
+ gfc_index_zero_node, GT_EXPR,
+ build_int_cst (gfc_array_index_type, -1),
+ gfc_finish_block (&loop_body));
+ /*
+ int volume = 1;
+ for (int i = 0; i < ndims - 2; i++)
+ {
+ int g = grains[i];
+ dims[i] = g / volume;
+ strides[i + 1] = fstrides[i + 1] / g;
+ volume = volume * dims[i];
+ }
+ */
+ tree volume = gfc_create_var (gfc_array_index_type, "volume");
+ gfc_add_modify (block, volume, gfc_index_one_node);
+
+ gfc_init_block (&loop_body);
+ tree grain = gfc_create_var (gfc_array_index_type, "grain");
+ gfc_add_modify (&loop_body, grain,
+ gfc_build_array_ref (grains, idx, NULL_TREE, true));
+ tree dims_i = gfc_build_array_ref (ompdims, idx, NULL_TREE, true);
+ gfc_add_modify (&loop_body, dims_i,
+ fold_build2_loc (input_location, TRUNC_DIV_EXPR,
+ gfc_array_index_type, grain, volume));
+ tree nidx = fold_build2_loc (input_location, PLUS_EXPR,
+ gfc_array_index_type, idx,
+ gfc_index_one_node);
+ tree strides_ni = gfc_build_array_ref (ompstrides, nidx, NULL_TREE,
true);
+ tree fstrides_ni = gfc_conv_descriptor_stride_get (desc, nidx);
+ gfc_add_modify (&loop_body, strides_ni,
+ fold_build2_loc (input_location, TRUNC_DIV_EXPR,
+ gfc_array_index_type, fstrides_ni,
+ grain));
+ gfc_add_modify (&loop_body, volume,
+ fold_build2_loc (input_location, MULT_EXPR,
+ gfc_array_index_type, volume, dims_i));
+
+ gfc_simple_for_loop (block, idx, gfc_index_zero_node,
+ build_int_cst (gfc_array_index_type, ndims - 2),
+ LT_EXPR, gfc_index_one_node,
+ gfc_finish_block (&loop_body));
+
+ /* dims[ndims - 2] = fstrides[ndims - 1] / volume; */
+ tree dimsm2
+ = gfc_build_array_ref (ompdims,
+ build_int_cst (gfc_array_index_type, ndims - 2),
+ NULL_TREE, true);
+ gfc_add_modify (block, dimsm2,
+ fold_build2_loc (input_location, TRUNC_DIV_EXPR,
+ gfc_array_index_type, fstridem1,
+ volume));
+ }
+}
+
+/* Return TRUE if update for N can definitely be done with a single contiguous
+ transfer. If no or if we can't tell, return FALSE. */
+
+static bool
+gfc_omp_contiguous_update_p (gfc_omp_namelist *n)
+{
+ gfc_expr *contig_expr = n->expr;
+
+ if (!n->expr)
+ {
+ if (n->sym->attr.contiguous)
+ return true;
+
+ tree desc = gfc_trans_omp_variable (n->sym, false);
+ tree type = TREE_TYPE (desc);
+ if (!GFC_ARRAY_TYPE_P (type) && !GFC_DESCRIPTOR_TYPE_P (type))
+ return true;
+
+ contig_expr = gfc_lval_expr_from_sym (n->sym);
+ }
+
+ return gfc_is_simply_contiguous (contig_expr, false, true);
+}
+
static tree
gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
locus where, bool declare_simd = false,
@@ -5400,6 +5744,184 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses *clauses,
default:
gcc_unreachable ();
}
+
+ if ((list == OMP_LIST_TO || list == OMP_LIST_FROM)
+ && (!n->expr || (n->expr->ref && n->expr->rank > 0))
+ && !gfc_omp_contiguous_update_p (n) && !iterator)
+ {
+ int ndims;
+ gfc_se se;
+ gfc_init_se (&se, NULL);
+
+ tree desc, span = NULL_TREE;
+
+ if (n->expr)
+ {
+ if (n->expr->rank)
+ gfc_conv_expr_descriptor (&se, n->expr);
+ else
+ gfc_conv_expr (&se, n->expr);
+
+ desc = se.expr;
+ /* The span is the distance between two array elements
+ along the innermost dimension (there may be padding
+ or other data between elements, e.g. of a derived-type
+ array). */
+ if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (desc)))
+ span = gfc_conv_descriptor_span_get (desc);
+ else
+ span = gfc_get_array_span (desc, n->expr);
+ ndims = n->expr->rank;
+ }
+ else
+ {
+ desc = gfc_trans_omp_variable (n->sym, false);
+ tree type = TREE_TYPE (desc);
+ if (GFC_DESCRIPTOR_TYPE_P (type))
+ span = gfc_conv_descriptor_span_get (desc);
+ ndims = GFC_TYPE_ARRAY_RANK (type);
+ }
+
+ gfc_add_block_to_block (&iter_block, &se.pre);
+
+ tree ompdims, ompstrides;
+
+ gfc_desc_to_omp_noncontig_array (&iter_block, &ompdims,
+ &ompstrides, desc, ndims);
+
+ tree type = TREE_TYPE (desc);
+ tree etype = gfc_get_element_type (type);
+ tree elsize = fold_convert (gfc_array_index_type,
+ size_in_bytes (etype));
+
+ tree ptr = gfc_conv_array_data (desc);
+ tree offset = gfc_conv_array_offset (desc);
+
+ if (!span)
+ /* The span is the element size. */
+ span = elsize;
+
+ tree node = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+
+ switch (list)
+ {
+ case OMP_LIST_TO:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO_GRID);
+ break;
+ case OMP_LIST_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FROM_GRID);
+ break;
+ default:
+ gcc_unreachable ();
+ }
+
+ gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
+ tree byte_offset = fold_convert (sizetype, offset);
+ byte_offset = size_binop (MULT_EXPR, byte_offset,
+ fold_convert (sizetype, span));
+ tree origin
+ = fold_build2_loc (input_location, POINTER_PLUS_EXPR,
+ TREE_TYPE (ptr), ptr, byte_offset);
+
+ OMP_CLAUSE_SIZE (node) = elsize;
+
+ omp_clauses = gfc_trans_add_clause (node, omp_clauses);
+
+ auto_vec<tree, 5> dims;
+
+ for (int r = 0; r < ndims; r++)
+ {
+ tree d
+ = gfc_build_array_ref (ompdims,
+ build_int_cst
+ (gfc_array_index_type, r),
+ NULL_TREE, true);
+ d = gfc_evaluate_now (d, &iter_block);
+ dims.safe_push (d);
+ }
+
+ for (int r = ndims - 1; r >= 0; r--)
+ {
+ tree stride_r, len_r, lowbound_r;
+
+ tree rcst = build_int_cst (gfc_array_index_type, r);
+
+ stride_r = gfc_build_array_ref (ompstrides, rcst,
+ NULL_TREE, true);
+ lowbound_r = gfc_conv_array_lbound (desc, r);
+ len_r
+ = fold_build2_loc (input_location, MINUS_EXPR,
+ gfc_array_index_type,
+ gfc_conv_array_ubound (desc, r),
+ lowbound_r);
+ len_r
+ = fold_build2_loc (input_location, PLUS_EXPR,
+ gfc_array_index_type,
+ len_r, gfc_index_one_node);
+
+ lowbound_r
+ = fold_build2_loc (input_location, MULT_EXPR,
+ gfc_array_index_type, lowbound_r,
+ stride_r);
+
+ stride_r = gfc_evaluate_now (stride_r, &iter_block);
+ lowbound_r = gfc_evaluate_now (lowbound_r, &iter_block);
+ len_r = gfc_evaluate_now (len_r, &iter_block);
+
+ tree dim = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (dim, GOMP_MAP_GRID_DIM);
+ OMP_CLAUSE_DECL (dim) = lowbound_r;
+ OMP_CLAUSE_SIZE (dim) = len_r;
+
+ omp_clauses = gfc_trans_add_clause (dim, omp_clauses);
+
+ if (!integer_onep (stride_r)
+ || (r == 0 && !operand_equal_p (span, elsize)))
+ {
+ tree snode = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (snode,
+ GOMP_MAP_GRID_STRIDE);
+ OMP_CLAUSE_DECL (snode) = stride_r;
+ if (r == 0 && !operand_equal_p (span, elsize))
+ OMP_CLAUSE_SIZE (snode) = span;
+ omp_clauses = gfc_trans_add_clause (snode,
+ omp_clauses);
+ }
+ }
+ origin = build_fold_indirect_ref (origin);
+ tree eltype = gfc_get_element_type (TREE_TYPE (desc));
+ tree arrtype
+ = gfc_trans_omp_arrayshape_type (eltype, &dims);
+ OMP_CLAUSE_DECL (node)
+ = build1_loc (input_location, VIEW_CONVERT_EXPR,
+ arrtype, origin);
+ if (!iterator)
+ gfc_add_block_to_block (block, &iter_block);
+ continue;
+ }
+
+ if (iterator && !gfc_omp_contiguous_update_p (n))
+ {
+ gfc_expr *contig_expr
+ = n->expr ? n->expr : gfc_lval_expr_from_sym (n->sym);
+ gfc_se cont_se;
+ gfc_init_se (&cont_se, NULL);
+ gfc_conv_is_contiguous_expr (&cont_se, contig_expr);
+ gfc_add_block_to_block (&iter_block, &cont_se.pre);
+ gfc_trans_runtime_check (
+ true, false,
+ fold_build1 (TRUTH_NOT_EXPR, boolean_type_node,
+ cont_se.expr),
+ &iter_block, &n->where,
+ "non-contiguous array section with iterator "
+ "in target update not supported");
+ gfc_add_block_to_block (&iter_block, &cont_se.post);
+ if (!n->expr)
+ gfc_free_expr (contig_expr);
+ }
+
tree node = build_omp_clause (gfc_get_location (&n->where),
clause_code);
if (n->expr == NULL
diff --git a/gcc/testsuite/gfortran.dg/gomp/bad-array-section-1.f90
b/gcc/testsuite/gfortran.dg/gomp/bad-array-section-1.f90
new file mode 100644
index 00000000000..58e47d927d7
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/bad-array-section-1.f90
@@ -0,0 +1,31 @@
+! { dg-do compile }
+!
+! A stride specified on an array section in a "target update" clause
+! must be a positive constant.
+
+module m
+ implicit none
+ integer :: a(10, 10)
+ integer :: b(10)
+contains
+ subroutine sub (n)
+ integer, value :: n
+
+ !$omp target update to(a(::-2, 2))
+! { dg-error "Stride must be positive in TO clause" "" { target *-*-* } .-1 }
+
+ !$omp target update to(a(10:1:-1, 2))
+! { dg-error "Stride must be positive in TO clause" "" { target *-*-* } .-1 }
+
+ !$omp target update from(a(::0, 2))
+! { dg-error "Illegal stride of zero" "" { target *-*-* } .-1 }
+! { dg-error "'a' in FROM clause at .1. is not a proper array section" "" {
target *-*-* } .-2 }
+
+ !$omp target update to(b(::-1))
+! { dg-error "Stride must be positive in TO clause" "" { target *-*-* } .-1 }
+
+ ! A non-constant stride cannot be checked at compile time; only a
+ ! runtime check can catch a negative value here.
+ !$omp target update to(b(1:10:n))
+ end subroutine sub
+end module m
diff --git a/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-1.f90
b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-1.f90
new file mode 100644
index 00000000000..5c60f5cac62
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-1.f90
@@ -0,0 +1,19 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+integer :: basicarray(100)
+integer, allocatable :: allocarray(:)
+
+allocate(allocarray(1:20))
+
+!$omp target update to(basicarray)
+
+!$omp target update from(basicarray(:))
+
+!$omp target update to(allocarray)
+
+!$omp target update from(allocarray(:))
+
+end
+
+! { dg-final { scan-tree-dump-times {omp target update from\(} 2 "original" } }
+! { dg-final { scan-tree-dump-times {omp target update to\(} 2 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-2.f90
b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-2.f90
new file mode 100644
index 00000000000..f5a52736b0c
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-2.f90
@@ -0,0 +1,16 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+integer, allocatable :: allocarray(:)
+integer, allocatable :: allocarray2(:,:)
+
+allocate(allocarray(1:20))
+allocate(allocarray2(1:20,1:20))
+
+! This one must be noncontiguous
+!$omp target update to(allocarray(::2))
+! { dg-final { scan-tree-dump {omp target update map\(to_grid:} "original" } }
+
+!$omp target update from(allocarray2(:,5:15))
+! { dg-final { scan-tree-dump {omp target update from\(} "original" } }
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-3.f90
b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-3.f90
new file mode 100644
index 00000000000..5cbfe7c7be5
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-3.f90
@@ -0,0 +1,16 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+integer, allocatable :: allocarray(:,:)
+
+allocate(allocarray(1:20,1:20))
+
+! This one could possibly be handled as a contiguous update - but isn't,
+! for now.
+!$omp target update to(allocarray(1:20,5:15))
+! { dg-final { scan-tree-dump {omp target update map\(to_grid:} "original" } }
+
+!$omp target update from(allocarray(:,5:15:2))
+! { dg-final { scan-tree-dump {omp target update map\(from_grid:} "original" }
}
+
+end
+
diff --git a/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-4.f90
b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-4.f90
new file mode 100644
index 00000000000..53152aacbb4
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-4.f90
@@ -0,0 +1,15 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+integer, target :: tgtarray(20)
+integer, pointer, contiguous :: arrayptr(:)
+
+arrayptr => tgtarray
+
+!$omp target update from(arrayptr)
+! { dg-final { scan-tree-dump {omp target update from\(} "original" } }
+
+!$omp target update to(arrayptr(::2))
+! { dg-final { scan-tree-dump {omp target update map\(to_grid:} "original" } }
+
+end
+
diff --git a/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-5.f90
b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-5.f90
new file mode 100644
index 00000000000..a964b59d268
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-5.f90
@@ -0,0 +1,28 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+! A derived-type component array accessed through the parent array is
+! noncontiguous, even with no explicit stride at the "target update" site --
the
+! noncontiguity is carried by the pointer's own runtime descriptor. OpenMP's
+! array-section grammar does not accept a bare "parent(:)%component" as a list
+! item directly, so this can only be reached via a POINTER.
+
+type t
+ complex(kind=8) :: c
+ integer :: i
+end type t
+
+type(t), target :: var(10)
+complex(kind=8), pointer :: ptr(:)
+
+! No explicit stride: intrinsically noncontiguous via the derived-type
+! component.
+ptr => var(:)%c
+!$omp target update from(ptr)
+! { dg-final { scan-tree-dump {omp target update map\(from_grid:} "original" }
}
+
+! Combined with an explicit stride on the pointer association.
+ptr => var(::2)%c
+!$omp target update to(ptr)
+! { dg-final { scan-tree-dump {omp target update map\(to_grid:} "original" } }
+
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-6.f90
b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-6.f90
new file mode 100644
index 00000000000..17dfb15c9f1
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-6.f90
@@ -0,0 +1,15 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+! The section has no explicit stride and the array is contiguous overall, but
+! for a genuinely runtime-variable bound the accessed memory may be
+! noncontiguous.
+
+subroutine s(n, m)
+ implicit none
+ integer :: n, m
+ integer :: array(n,n)
+
+ !$omp target update to(array(1:m,:))
+ ! { dg-final { scan-tree-dump {omp target update map\(to_grid:} "original" }
}
+
+end subroutine s
diff --git a/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-7.f90
b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-7.f90
new file mode 100644
index 00000000000..92dfa2927d3
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-7.f90
@@ -0,0 +1,17 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original" }
+
+! arr(1:m,:,k) has three subscripts written (u.ar.dimen == 3), but the fixed
+! index k collapses one dimension, so the section's rank is only 2.
+! Confirm exactly two grid_dim clauses are emitted, matching the section's
+! rank, not three.
+
+subroutine s(arr, n, m, k)
+ implicit none
+ integer :: n, m, k
+ integer :: arr(n, n, n)
+
+ !$omp target update from(arr(1:m,:,k))
+ ! { dg-final { scan-tree-dump-times {map\(grid_dim:} 2 "original" } }
+
+end subroutine s
diff --git a/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-8.f90
b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-8.f90
new file mode 100644
index 00000000000..9343f68352e
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-8.f90
@@ -0,0 +1,37 @@
+! { dg-do compile }
+
+! Strided array sections in "target update" that are not a plain array
+! section at the end of a component chain are rejected.
+
+program p
+implicit none
+
+type u
+ integer :: lo, b, hi
+end type
+
+type t2
+ type(u) :: a(10)
+end type
+
+type tc
+ character(len=4) :: c(10)
+end type
+
+type(t2) :: y
+type(tc) :: z
+character(len=4) :: s(10)
+
+! Strided section in the middle of the component chain.
+!$omp target update from(y%a(1:10:2)%b)
+! { dg-error "'y' in FROM clause at .1. is not a proper array section" "" {
target *-*-* } .-1 }
+
+! Substring as the last reference of a strided section.
+!$omp target update from(z%c(1:10:2)(2:3))
+! { dg-error "'z' in FROM clause at .1. is not a proper array section" "" {
target *-*-* } .-1 }
+
+! Substring of an array element.
+!$omp target update from(s(3)(2:3))
+! { dg-error "Unexpected substring reference in FROM clause" "" { target *-*-*
} .-1 }
+
+end program p
diff --git a/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-9.f90
b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-9.f90
new file mode 100644
index 00000000000..08bc62f90e6
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/noncontig-updates-9.f90
@@ -0,0 +1,31 @@
+! { dg-do compile }
+
+! The per-dimension checks on a "target update" array section must be applied
+! to every dimension, not only up to the first one with a constant stride.
+
+subroutine s (a, b)
+ implicit none
+ integer :: a(10,10), b(10,10,10)
+
+ ! Valid: constant positive strides in several dimensions.
+ !$omp target update to(a(1:10:2, 1:10:3))
+ !$omp target update to(b(1:10:2, 2:8, 1:10:5))
+
+ ! A bad second dimension after a valid constant stride in the first.
+ !$omp target update to(a(1:10:2, 10:1:-1))
+ ! { dg-error "Stride must be positive in TO clause at .1." "" { target *-*-*
} .-1 }
+ !$omp target update from(a(1:10:2, [1,3,5]))
+ ! { dg-error "'a' in FROM clause at .1. is not a proper array section" "" {
target *-*-* } .-1 }
+
+ ! Same, but the bad dimension is the third one.
+ !$omp target update to(b(1:10:2, 2:8, 10:1:-2))
+ ! { dg-error "Stride must be positive in TO clause at .1." "" { target *-*-*
} .-1 }
+ !$omp target update from(b(1:10:2, 1:10:3, [1,3,5]))
+ ! { dg-error "'b' in FROM clause at .1. is not a proper array section" "" {
target *-*-* } .-1 }
+
+ ! Controls: the bad dimension comes first.
+ !$omp target update to(a(10:1:-1, 1:10:2))
+ ! { dg-error "Stride must be positive in TO clause at .1." "" { target *-*-*
} .-1 }
+ !$omp target update from(a([1,3,5], 1:10:2))
+ ! { dg-error "'a' in FROM clause at .1. is not a proper array section" "" {
target *-*-* } .-1 }
+end subroutine
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90
b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90
index d9c92cf4679..f7601053ea7 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90
@@ -17,7 +17,14 @@ program test
!$omp target update from (iterator(i=1:DIM1): z(i)%ptr(:))
end program
+! Since these array sections are over POINTER components that cannot be proven
+! contiguous at compile time, and since the map iterators with strided updates
+! are not yet supported, these now always take the plain (non-grid) path with a
+! runtime contiguity check instead.
! { dg-final { scan-tree-dump-times "if \\(i <= 17\\) goto <D\\\.\[0-9\]+>;
else goto <D\\\.\[0-9\]+>;" 2 "gimple" } }
! { dg-final { scan-tree-dump-times "if \\(j <= 39\\) goto <D\\\.\[0-9\]+>;
else goto <D\\\.\[0-9\]+>;" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "to\\(iterator\\(integer\\(kind=4\\)
j=1:39:1, integer\\(kind=4\\) i=1:17:1, loop_label=<D\\\.\[0-9\]+>,
elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):MEM <\[^>\]+>
\\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]" 2 "gimple" } }
-! { dg-final { scan-tree-dump-times "from\\(iterator\\(integer\\(kind=4\\)
i=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+,
index=D\\\.\[0-9\]+\\):MEM <\[^>\]+> \\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]"
1 "gimple" } }
+! { dg-final { scan-tree-dump-times "non-contiguous array section with
iterator in target update not supported" 3 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target update
to\\(iterator\\(integer\\(kind=4\\) j=1:39:1, integer\\(kind=4\\) i=1:17:1,
loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+,
index=D\\\.\[0-9\]+\\):" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target update
from\\(iterator\\(integer\\(kind=4\\) i=1:17:1, loop_label=<D\\\.\[0-9\]+>,
elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):" 1 "gimple" } }
+! { dg-final { scan-tree-dump-not "to_grid:" "gimple" } }
+! { dg-final { scan-tree-dump-not "from_grid:" "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-4.f90
b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-4.f90
new file mode 100644
index 00000000000..f4566e0445f
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-4.f90
@@ -0,0 +1,21 @@
+! { dg-do compile }
+! { dg-options "-fopenmp" }
+
+! A fully-indexed element reference (all subscripts scalar, no ':') into a
+! bare POINTER array (no CONTIGUOUS attribute) is trivially contiguous - it's
+! a single scalar value - regardless of the pointer's own contiguity. This
+! must not be routed through the noncontiguous/"grid" update lowering, which
+! assumes an array-section (rank >= 1) expression throughout and ICEs in
+! gfc_get_element_type otherwise.
+
+program test
+ implicit none
+
+ integer, parameter :: DIM1 = 10
+
+ integer, pointer :: arr(:,:)
+
+ allocate (arr(DIM1, DIM1))
+
+ !$omp target update to (iterator(i=1:DIM1): arr(i,i))
+end program
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-5.f90
b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-5.f90
new file mode 100644
index 00000000000..3aabcbdf8cf
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-5.f90
@@ -0,0 +1,23 @@
+! { dg-do compile }
+! { dg-options "-fopenmp" }
+
+! Check restrictions on non-contiguous arrays and map iterators.
+
+program test
+ implicit none
+ integer :: array2(10, 10)
+ integer :: y(9, 10)
+
+ ! Explicit non-unit literal stride + iterator
+ !$omp target update to(iterator(i=1:9): array2(::2, i)) ! { dg-error
"Non-contiguous array with map iterator" }
+ ! { dg-error "Non-unit stride with iterator in TO clause" "" { target *-*-*
} .-1 }
+
+ ! A partial range in the first dimension followed by a fixed index is still
+ ! contiguous. Confirms the contiguity check does not over-trigger on such a
+ ! shape.
+ !$omp target update to(iterator(i=1:9): array2(1:5, i))
+
+ ! Iterator as the first dim, with a section of the outer dim -- a
+ ! strided/noncontiguous access.
+ !$omp target update to(iterator(i=1:9): y(i, 1:5)) ! { dg-error
"Non-contiguous array with map iterator" }
+end program
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-6.f90
b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-6.f90
new file mode 100644
index 00000000000..5a3ab55d607
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-6.f90
@@ -0,0 +1,39 @@
+! { dg-do compile }
+! { dg-options "-fopenmp -fdump-tree-gimple" }
+
+! Compile-time (dump-scan) companion to the real-data execution test in
+! libgomp.fortran/target-update-iterators-5.f90, for the same bug: a PLAIN
+! (non-pointer-component) Fortran array combined with a map iterator over
+! whole array sections -- Tobias Burnus's own "should be fine" example,
+! y(:,:,i,2) -- silently loses the iterator entirely.
+!
+! Confirmed here at the gimple level: unlike target-update-iterators-3.f90's
+! POINTER-component case (which correctly keeps the loop/"iterator(...)"
+! structure -- see the loop_label=/elems=omp_iter_data patterns there), this
+! plain-array case produces no loop, no omp_iter_data, and no
+! "iterator(...)" on the pragma at all. The whole update collapses to a
+! single, unconditional, non-iterated "target update from" statement.
+!
+! Root cause: remove_unused_omp_iterator_vars (gimplify.cc:9970) walks each
+! clause's DECL/SIZE looking for a literal reference to the iterator
+! variable; gfc_trans_omp_clauses evaluates a plain array section's
+! address/length into scalar temporaries (gfc_evaluate_now) that no longer
+! contain that literal reference, so the walk finds nothing and the
+! iterator is dropped as apparently unused. Not yet fixed.
+
+program test
+ implicit none
+ integer, parameter :: D1 = 4, D2 = 4, D3 = 9, D4 = 3
+ integer :: y(D1, D2, D3, D4)
+
+ !$omp target update from(iterator(k=1:D3): y(:, :, k, 2))
+end program
+
+! The iterator has been silently dropped: no loop, no per-iteration data
+! capture, and no "iterator(...)" on the pragma -- contrast with
+! target-update-iterators-3.f90's pointer-component case, which keeps all
+! three. Exactly one (non-iterated) "target update from" is emitted.
+! { dg-final { scan-tree-dump-not "omp_iter_data" "gimple" } }
+! { dg-final { scan-tree-dump-not "loop_label=" "gimple" } }
+! { dg-final { scan-tree-dump-not "iterator\\(" "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target update from" 1
"gimple" } }
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-1.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-1.f90
new file mode 100644
index 00000000000..6ee87e8043b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-1.f90
@@ -0,0 +1,54 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+implicit none
+integer, allocatable, target :: arr(:), arr2(:,:)
+integer, pointer :: ap(:), ap2(:,:)
+integer :: i, j
+
+allocate(arr(1:20))
+
+arr = 0
+
+!$omp target enter data map(to: arr)
+
+ap => arr(1:20:2)
+ap = 5
+
+!$omp target update to(ap)
+
+!$omp target exit data map(from: arr)
+
+do i=1,20
+ if (mod(i,2).eq.1.and.arr(i).ne.5) stop 1
+ if (mod(i,2).eq.0.and.arr(i).ne.0) stop 2
+end do
+
+allocate(arr2(1:20,1:20))
+
+ap2 => arr2(2:10:2,3:12:3)
+
+arr2 = 1
+
+!$omp target enter data map(to: arr2)
+
+!$omp target
+ap2 = 5
+!$omp end target
+
+!$omp target update from(ap2)
+
+do i=1,20
+ do j=1,20
+ if (i.ge.2.and.i.le.10.and.mod(i-2,2).eq.0.and.&
+ &j.ge.3.and.j.le.12.and.mod(j-3,3).eq.0) then
+ if (arr2(i,j).ne.5) stop 3
+ else
+ if (arr2(i,j).ne.1) stop 4
+ end if
+ end do
+end do
+
+!$omp target exit data map(delete: arr2)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-10.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-10.f90
new file mode 100644
index 00000000000..c47ce38918d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-10.f90
@@ -0,0 +1,29 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+character(len=8), allocatable, dimension(:) :: lines
+integer :: i
+
+allocate(lines(10))
+
+lines = "OMPHELLO"
+
+!$omp target enter data map(to: lines)
+
+!$omp target
+lines = "NEWVALUE"
+!$omp end target
+
+!$omp target update from(lines(5:7:2))
+
+do i=1,10
+ if (i.eq.5.or.i.eq.7) then
+ if (lines(i).ne."NEWVALUE") stop 1
+ else
+ if (lines(i).ne."OMPHELLO") stop 2
+ end if
+end do
+
+!$omp target exit data map(delete: lines)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-11.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-11.f90
new file mode 100644
index 00000000000..a93acf21d77
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-11.f90
@@ -0,0 +1,51 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+program p
+implicit none
+real(kind=4) :: arr(10,10,10,10)
+
+call s(arr,9,9,9,9)
+
+contains
+
+subroutine s(arr,m,n,o,p)
+implicit none
+integer :: i,m,n,o,p
+integer :: a,b,c,d
+real(kind=4) :: arr(0:m,0:n,0:o,0:p)
+
+arr = 0
+
+!$omp target enter data map(to: arr)
+
+!$omp target
+do i=0,9
+ arr(i,i,i,i) = i
+end do
+!$omp end target
+
+!$omp target update from(arr(0:2,0:2,0:2,0:2))
+
+do a=0,9
+ do b=0,9
+ do c=0,9
+ do d=0,9
+ if (a.le.2.and.b.le.2.and.c.le.2.and.d.le.2) then
+ if (a.eq.b.and.b.eq.c.and.c.eq.d) then
+ if (arr(a,b,c,d).ne.a) stop 1
+ else
+ if (arr(a,b,c,d).ne.0) stop 2
+ end if
+ else
+ if (arr(a,b,c,d).ne.0) stop 3
+ end if
+ end do
+ end do
+ end do
+end do
+
+!$omp target exit data map(delete: arr)
+
+end subroutine s
+end program p
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-12.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-12.f90
new file mode 100644
index 00000000000..c47fbdb0d11
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-12.f90
@@ -0,0 +1,59 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! Test plain, fixed-size arrays, and also pointers to same.
+
+implicit none
+integer(kind=8) :: arr(10,30)
+integer, target :: arr2(9,11,13)
+integer, pointer :: parr(:,:,:)
+integer :: i, j, k
+
+arr = 0
+!$omp target enter data map(to: arr)
+
+!$omp target
+arr = 99
+!$omp end target
+
+!$omp target update from(arr(1:10:3,5:30:7))
+
+do i=1,10
+ do j=1,30
+ if (mod(i-1,3).eq.0.and.mod(j-5,7).eq.0) then
+ if (arr(i,j).ne.99) stop 1
+ else
+ if (arr(i,j).ne.0) stop 2
+ endif
+ end do
+end do
+
+!$omp target exit data map(delete: arr)
+
+arr2 = 0
+parr => arr2
+!$omp target enter data map(to: parr)
+
+!$omp target
+parr = 99
+!$omp end target
+
+!$omp target update from(parr(7:9:2,5:7:2,3:6:3))
+
+do i=1,9
+ do j=1,11
+ do k=1,13
+ if (i.ge.7.and.j.ge.5.and.k.ge.3.and.&
+ &i.le.9.and.j.le.7.and.k.le.6.and.&
+ &mod(i-7,2).eq.0.and.mod(j-5,2).eq.0.and.mod(k-3,3).eq.0) then
+ if (parr(i,j,k).ne.99) stop 3
+ else
+ if (parr(i,j,k).ne.0) stop 4
+ end if
+ end do
+ end do
+end do
+
+!$omp target exit data map(delete: parr)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-13.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-13.f90
new file mode 100644
index 00000000000..42f867efefc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-13.f90
@@ -0,0 +1,42 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+implicit none
+integer, allocatable :: arr(:,:,:,:,:)
+integer :: i, j, k, l, m
+
+allocate (arr(18,19,20,21,22))
+
+arr = 0
+
+!$omp target enter data map(to: arr)
+
+arr = 10
+
+!$omp target update to(arr(1:3:2,1:4:3,1:5:4,1:6:5,1:7:6))
+
+!$omp target
+do i=1,18
+ do j=1,19
+ do k=1,20
+ do l=1,21
+ do m=1,22
+ if ((i.eq.1.or.i.eq.3).and.&
+ &(j.eq.1.or.j.eq.4).and.&
+ &(k.eq.1.or.k.eq.5).and.&
+ &(l.eq.1.or.l.eq.6).and.&
+ &(m.eq.1.or.m.eq.7)) then
+ if (arr(i,j,k,l,m).ne.10) stop 1
+ else
+ if (arr(i,j,k,l,m).ne.0) stop 2
+ end if
+ end do
+ end do
+ end do
+ end do
+end do
+!$omp end target
+
+!$omp target exit data map(delete: arr)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-14.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-14.f90
new file mode 100644
index 00000000000..b5b8a0e1425
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-14.f90
@@ -0,0 +1,41 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! Check that the size of the array section to update is properly computed when
+! the last (i.e. outer) dimension is a fixed index.
+
+program p
+implicit none
+integer :: arr(10,5,9)
+integer :: i,j,k,m
+
+m = 6
+arr = -1
+
+!$omp target enter data map(to: arr)
+
+!$omp target
+do k=1,9
+ do j=1,5
+ do i=1,10
+ arr(i,j,k) = i*1000 + j*10 + k
+ end do
+ end do
+end do
+!$omp end target
+
+!$omp target update from(arr(1:m,:,3))
+
+do j=1,5
+ do i=1,10
+ if (i.le.m) then
+ if (arr(i,j,3).ne.i*1000+j*10+3) stop 1
+ else
+ if (arr(i,j,3).ne.-1) stop 2
+ end if
+ end do
+end do
+
+!$omp target exit data map(delete: arr)
+
+end program p
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-15.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-15.f90
new file mode 100644
index 00000000000..e655fb0976d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-15.f90
@@ -0,0 +1,95 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! Check that a strided array section reached through a derived-type component
+! is updated as a non-contiguous section, i.e. only the selected elements are
+! transferred. Cover plain, pointer and allocatable array components, and a
+! whole pointer component whose target is itself strided.
+
+program p
+implicit none
+
+type t
+ integer :: arr(10)
+end type
+
+type tp
+ integer, pointer :: p(:)
+ integer, allocatable :: al(:)
+end type
+
+type(t) :: x
+type(tp) :: w
+integer, target :: big(20)
+integer :: i
+
+big = [(i, i = 1, 20)]
+x%arr = [(i, i = 1, 10)]
+
+!$omp target enter data map(to: x)
+
+!$omp target
+ x%arr = -x%arr
+!$omp end target
+
+!$omp target update from(x%arr(1:10:2))
+
+do i = 1, 10
+ if (mod (i, 2) == 1) then
+ if (x%arr(i) /= -i) stop 1
+ else
+ if (x%arr(i) /= i) stop 2
+ end if
+end do
+
+!$omp target exit data map(delete: x)
+
+allocate (w%p(10), w%al(10))
+w%p = [(i, i = 1, 10)]
+w%al = [(i, i = 1, 10)]
+
+!$omp target enter data map(to: w, w%p, w%al)
+
+!$omp target
+ w%p = -w%p
+ w%al = -w%al
+!$omp end target
+
+!$omp target update from(w%p(1:10:2))
+!$omp target update from(w%al(1:10:2))
+
+do i = 1, 10
+ if (mod (i, 2) == 1) then
+ if (w%p(i) /= -i) stop 3
+ if (w%al(i) /= -i) stop 4
+ else
+ if (w%p(i) /= i) stop 5
+ if (w%al(i) /= i) stop 6
+ end if
+end do
+
+!$omp target exit data map(delete: w%p, w%al, w)
+
+! A whole pointer component (no section) aimed at a strided target.
+deallocate (w%p, w%al)
+w%p => big(1:20:2)
+
+!$omp target enter data map(to: big)
+
+!$omp target
+ big = -big
+!$omp end target
+
+!$omp target update from(w%p)
+
+do i = 1, 20
+ if (mod (i, 2) == 1) then
+ if (big(i) /= -i) stop 7
+ else
+ if (big(i) /= i) stop 8
+ end if
+end do
+
+!$omp target exit data map(delete: big)
+
+end program p
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-16.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-16.f90
new file mode 100644
index 00000000000..0fbcf99ab62
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-16.f90
@@ -0,0 +1,163 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! Check that a whole multi-dimensional array, or a section of it, reached
+! through a pointer to a derived-type component is updated correctly. The
+! descriptor's span (the size of the derived type) then differs from the
+! element size, which must be taken into account for the outer dimensions as
+! well as the innermost one.
+
+program noncontig_updates_16
+implicit none
+
+type t
+ integer :: x, y
+end type
+
+type(t), target :: a(10,10)
+type(t), target :: b(6,5,4)
+integer, pointer :: p(:,:)
+integer, pointer :: q(:,:,:)
+integer :: i, j, k
+
+! 2-D.
+do j = 1, 10
+ do i = 1, 10
+ a(i,j)%x = i + 100*j
+ a(i,j)%y = -1
+ end do
+end do
+p => a%x
+
+!$omp target enter data map(to: a)
+
+!$omp target
+ do j = 1, 10
+ do i = 1, 10
+ a(i,j)%x = a(i,j)%x + 10000
+ a(i,j)%y = -2
+ end do
+ end do
+!$omp end target
+
+!$omp target update from(p)
+
+! Only the 'x' components must have been updated.
+do j = 1, 10
+ do i = 1, 10
+ if (a(i,j)%x /= i + 100*j + 10000) stop 1
+ if (a(i,j)%y /= -1) stop 2
+ end do
+end do
+
+! 2-D section: only the 'x' components of a(1:5,1:5) must be updated.
+do j = 1, 10
+ do i = 1, 10
+ a(i,j)%x = i + 100*j
+ a(i,j)%y = -1
+ end do
+end do
+
+!$omp target update to(a)
+
+!$omp target
+ do j = 1, 10
+ do i = 1, 10
+ a(i,j)%x = a(i,j)%x + 10000
+ a(i,j)%y = -2
+ end do
+ end do
+!$omp end target
+
+!$omp target update from(p(1:5,1:5))
+
+do j = 1, 10
+ do i = 1, 10
+ if (i <= 5 .and. j <= 5) then
+ if (a(i,j)%x /= i + 100*j + 10000) stop 5
+ else
+ if (a(i,j)%x /= i + 100*j) stop 6
+ end if
+ if (a(i,j)%y /= -1) stop 7
+ end do
+end do
+
+!$omp target exit data map(delete: a)
+
+! 3-D.
+do k = 1, 4
+ do j = 1, 5
+ do i = 1, 6
+ b(i,j,k)%x = i + 10*j + 100*k
+ b(i,j,k)%y = -1
+ end do
+ end do
+end do
+q => b%x
+
+!$omp target enter data map(to: b)
+
+!$omp target
+ do k = 1, 4
+ do j = 1, 5
+ do i = 1, 6
+ b(i,j,k)%x = b(i,j,k)%x + 10000
+ b(i,j,k)%y = -2
+ end do
+ end do
+ end do
+!$omp end target
+
+!$omp target update from(q)
+
+do k = 1, 4
+ do j = 1, 5
+ do i = 1, 6
+ if (b(i,j,k)%x /= i + 10*j + 100*k + 10000) stop 3
+ if (b(i,j,k)%y /= -1) stop 4
+ end do
+ end do
+end do
+
+! 3-D section: only the 'x' components of b(2:5,2:4,2:3) must be updated.
+do k = 1, 4
+ do j = 1, 5
+ do i = 1, 6
+ b(i,j,k)%x = i + 10*j + 100*k
+ b(i,j,k)%y = -1
+ end do
+ end do
+end do
+
+!$omp target update to(b)
+
+!$omp target
+ do k = 1, 4
+ do j = 1, 5
+ do i = 1, 6
+ b(i,j,k)%x = b(i,j,k)%x + 10000
+ b(i,j,k)%y = -2
+ end do
+ end do
+ end do
+!$omp end target
+
+!$omp target update from(q(2:5,2:4,2:3))
+
+do k = 1, 4
+ do j = 1, 5
+ do i = 1, 6
+ if (i >= 2 .and. i <= 5 .and. j >= 2 .and. j <= 4 &
+ .and. k >= 2 .and. k <= 3) then
+ if (b(i,j,k)%x /= i + 10*j + 100*k + 10000) stop 8
+ else
+ if (b(i,j,k)%x /= i + 10*j + 100*k) stop 9
+ end if
+ if (b(i,j,k)%y /= -1) stop 10
+ end do
+ end do
+end do
+
+!$omp target exit data map(delete: b)
+
+end program noncontig_updates_16
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-17.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-17.f90
new file mode 100644
index 00000000000..eafb8f7f61a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-17.f90
@@ -0,0 +1,36 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! Check that a strided array section with a non-unit lower bound is looked up
+! at the right place in the device mapping. The lookup bias must be based on
+! the lower bound alone, not multiplied by the stride a second time.
+
+program noncontig_updates_17
+implicit none
+
+integer :: a(100), i
+
+a = [(i, i = 1, 100)]
+
+!$omp target enter data map(to: a(31:61))
+
+!$omp target
+ do i = 31, 61
+ a(i) = -i
+ end do
+!$omp end target
+
+! Elements 31, 41, 51 and 61.
+!$omp target update from(a(31:61:10))
+
+do i = 1, 100
+ if (i >= 31 .and. i <= 61 .and. mod (i - 31, 10) == 0) then
+ if (a(i) /= -i) stop 1
+ else
+ if (a(i) /= i) stop 2
+ end if
+end do
+
+!$omp target exit data map(delete: a(31:61))
+
+end program noncontig_updates_17
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-18.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-18.f90
new file mode 100644
index 00000000000..1ef1e2bcbfa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-18.f90
@@ -0,0 +1,150 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! Check sections of a pointer to a derived-type component, whose descriptor
+! span (the size of the derived type) differs from the element size: a
+! rank-1 section, a strided section, the 'to' direction, and an allocatable
+! array whose pointer has non-default lower bounds. Only the 'x' components
+! within the section may be transferred; the 'y' components and everything
+! outside the section must keep their values.
+
+program noncontig_updates_18
+implicit none
+type t
+ integer :: x, y
+end type
+type(t), target :: a(10,10)
+type(t), allocatable, target :: c(:,:)
+integer, pointer :: p(:,:), r(:,:)
+integer :: i, j, err
+logical :: in
+
+allocate (c(10,10))
+p => a%x
+r(0:,0:) => c%x
+
+! Rank-1 section of a rank-2 pointer.
+do j = 1, 10
+ do i = 1, 10
+ a(i,j)%x = i + 100*j
+ a(i,j)%y = -1
+ end do
+end do
+!$omp target enter data map(to: a)
+!$omp target
+ do j = 1, 10
+ do i = 1, 10
+ a(i,j)%x = a(i,j)%x + 10000
+ a(i,j)%y = -2
+ end do
+ end do
+!$omp end target
+!$omp target update from(p(1:5,3))
+do j = 1, 10
+ do i = 1, 10
+ in = i <= 5 .and. j == 3
+ if (in) then
+ if (a(i,j)%x /= i + 100*j + 10000) stop 11
+ else
+ if (a(i,j)%x /= i + 100*j) stop 12
+ end if
+ if (a(i,j)%y /= -1) stop 13
+ end do
+end do
+!$omp target exit data map(delete: a)
+
+! Strided section.
+do j = 1, 10
+ do i = 1, 10
+ a(i,j)%x = i + 100*j
+ a(i,j)%y = -1
+ end do
+end do
+!$omp target enter data map(to: a)
+!$omp target
+ do j = 1, 10
+ do i = 1, 10
+ a(i,j)%x = a(i,j)%x + 10000
+ a(i,j)%y = -2
+ end do
+ end do
+!$omp end target
+!$omp target update from(p(1:9:2, 2:8:3))
+do j = 1, 10
+ do i = 1, 10
+ in = mod (i-1, 2) == 0 .and. i <= 9 .and. j >= 2 .and. j <= 8 .and. mod
(j-2, 3) == 0
+ if (in) then
+ if (a(i,j)%x /= i + 100*j + 10000) stop 21
+ else
+ if (a(i,j)%x /= i + 100*j) stop 22
+ end if
+ if (a(i,j)%y /= -1) stop 23
+ end do
+end do
+!$omp target exit data map(delete: a)
+
+! 'to' direction: the host changes everything, but only the section must
+! arrive on the device.
+do j = 1, 10
+ do i = 1, 10
+ a(i,j)%x = i + 100*j
+ a(i,j)%y = -1
+ end do
+end do
+!$omp target enter data map(to: a)
+do j = 1, 10
+ do i = 1, 10
+ a(i,j)%x = -5
+ a(i,j)%y = -6
+ end do
+end do
+!$omp target update to(p(2:4, 3:6))
+err = 0
+!$omp target map(tofrom: err)
+ do j = 1, 10
+ do i = 1, 10
+ in = i >= 2 .and. i <= 4 .and. j >= 3 .and. j <= 6
+ if (in) then
+ if (a(i,j)%x /= -5) err = err + 1
+ else
+ if (a(i,j)%x /= i + 100*j) err = err + 1
+ end if
+ if (a(i,j)%y /= -1) err = err + 1
+ end do
+ end do
+!$omp end target
+if (err /= 0) then
+ stop 31
+end if
+!$omp target exit data map(delete: a)
+
+! Allocatable array, pointer with lower bounds 0.
+do j = 1, 10
+ do i = 1, 10
+ c(i,j)%x = i + 100*j
+ c(i,j)%y = -1
+ end do
+end do
+!$omp target enter data map(to: c)
+!$omp target
+ do j = 1, 10
+ do i = 1, 10
+ c(i,j)%x = c(i,j)%x + 10000
+ c(i,j)%y = -2
+ end do
+ end do
+!$omp end target
+!$omp target update from(r(2:4, 3:6))
+do j = 1, 10
+ do i = 1, 10
+ in = i >= 3 .and. i <= 5 .and. j >= 4 .and. j <= 7
+ if (in) then
+ if (c(i,j)%x /= i + 100*j + 10000) stop 41
+ else
+ if (c(i,j)%x /= i + 100*j) stop 42
+ end if
+ if (c(i,j)%y /= -1) stop 43
+ end do
+end do
+!$omp target exit data map(delete: c)
+end program noncontig_updates_18
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-2.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-2.f90
new file mode 100644
index 00000000000..2d3efb8bfcc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-2.f90
@@ -0,0 +1,101 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+program p
+implicit none
+integer, allocatable, target :: arr3(:,:,:)
+integer, pointer :: ap3(:,:,:)
+integer :: i, j, k
+
+allocate(arr3(1:10,1:10,1:10))
+
+! CHECK 1
+
+arr3 = 0
+ap3 => arr3(1:10,1:10,1:10:2)
+
+!$omp target enter data map(to: arr3)
+
+!$omp target
+ap3 = 5
+!$omp end target
+
+!$omp target update from(ap3)
+
+call check(arr3, 0, 1, 1, 2)
+
+!$omp target exit data map(delete: arr3)
+
+! CHECK 2
+
+arr3 = 0
+ap3 => arr3(1:10,1:10:2,1:10)
+
+!$omp target enter data map(to: arr3)
+
+!$omp target
+ap3 = 5
+!$omp end target
+
+!$omp target update from(ap3)
+
+call check(arr3, 2, 1, 2, 1)
+
+!$omp target exit data map(delete: arr3)
+
+! CHECK 3
+
+arr3 = 0
+ap3 => arr3(1:10:2,1:10,1:10)
+
+!$omp target enter data map(to: arr3)
+
+!$omp target
+ap3 = 5
+!$omp end target
+
+!$omp target update from(ap3)
+
+call check(arr3, 4, 2, 1, 1)
+
+!$omp target exit data map(delete: arr3)
+
+! CHECK 4
+
+arr3 = 0
+ap3 => arr3(1:10:2,1:10:2,1:10:2)
+
+!$omp target enter data map(to: arr3)
+
+!$omp target
+ap3 = 5
+!$omp end target
+
+!$omp target update from(ap3)
+
+call check(arr3, 6, 2, 2, 2)
+
+!$omp target exit data map(delete: arr3)
+
+contains
+
+subroutine check(arr,cb,s1,s2,s3)
+implicit none
+integer :: arr(:,:,:)
+integer :: cb, s1, s2, s3
+
+do i=1,10
+ do j=1,10
+ do k=1,10
+ if (mod(k-1,s1).eq.0.and.mod(j-1,s2).eq.0.and.mod(i-1,s3).eq.0) then
+ if (arr(k,j,i).ne.5) stop cb+1
+ else
+ if (arr(k,j,i).ne.0) stop cb+2
+ end if
+ end do
+ end do
+end do
+
+end subroutine check
+
+end program p
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-3.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-3.f90
new file mode 100644
index 00000000000..14f1288a697
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-3.f90
@@ -0,0 +1,47 @@
+program p
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+integer :: A(200)
+A = [(i, i=1,200)]
+!$omp target enter data map(to: A(40:200))
+call foo(A(101:))
+
+contains
+
+subroutine foo(x)
+integer, target :: x(100)
+integer, pointer :: p(:,:)
+integer :: i, j
+
+p(0:5,-5:-1) => x(::2)
+
+!$omp target
+x = x * 2
+!$omp end target
+
+!$omp target update from(x(1:20:2))
+
+do i=1,20
+if (mod(i,2).eq.1 .and. x(i).ne.(100+i)*2) stop 1
+if (mod(i,2).eq.0 .and. x(i).ne.100+i) stop 2
+end do
+
+!$omp target
+p = 0
+!$omp end target
+
+!$omp target update from(p(::3,::2))
+
+do i=0,5
+ do j=-5,-1
+ if (mod(i,3).eq.0 .and. mod(j+5,2).eq.0) then
+ if (p(i,j).ne.0) stop 3
+ else
+ if (p(i,j).eq.0) stop 4
+ end if
+ end do
+end do
+
+end subroutine foo
+end program p
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-4.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-4.f90
new file mode 100644
index 00000000000..46e8c23d285
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-4.f90
@@ -0,0 +1,78 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+type t
+ complex(kind=8) :: c
+ integer :: i
+end type t
+
+type u
+ integer :: i, j
+ complex(kind=8) :: c
+ integer :: k
+end type u
+
+type(t), target :: var(10)
+type(u), target :: var2(10)
+complex(kind=8), pointer :: ptr(:)
+integer :: i
+
+do i=1,10
+ var(i)%c = dcmplx(i,0)
+ var(i)%i = i
+end do
+
+ptr => var(:)%c
+
+!$omp target enter data map(to: var)
+
+!$omp target
+var(:)%c = dcmplx(0,0)
+var(:)%i = 0
+!$omp end target
+
+!$omp target update from(ptr)
+
+do i=1,10
+ if (var(i)%c.ne.dcmplx(0,0)) stop 1
+ if (var(i)%i.ne.i) stop 2
+end do
+
+!$omp target exit data map(delete: var)
+
+! Now do it again with a differently-ordered derived type.
+
+do i=1,10
+ var2(i)%c = dcmplx(0,i)
+ var2(i)%i = i
+ var2(i)%j = i * 2
+ var2(i)%k = i * 3
+end do
+
+ptr => var2(::2)%c
+
+!$omp target enter data map(to: var2)
+
+!$omp target
+var2(:)%c = dcmplx(0,0)
+var2(:)%i = 0
+var2(:)%j = 0
+var2(:)%k = 0
+!$omp end target
+
+!$omp target update from(ptr)
+
+do i=1,10
+ if (mod(i,2).eq.1) then
+ if (var2(i)%c.ne.dcmplx(0,0)) stop 3
+ else
+ if (var2(i)%c.ne.dcmplx(0,i)) stop 4
+ end if
+ if (var2(i)%i.ne.i) stop 5
+ if (var2(i)%j.ne.i * 2) stop 6
+ if (var2(i)%k.ne.i * 3) stop 7
+end do
+
+!$omp target exit data map(delete: var2)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-5.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-5.f90
new file mode 100644
index 00000000000..9cc20fa321e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-5.f90
@@ -0,0 +1,55 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! Only some of an array mapped on the target
+
+integer, target :: arr(100)
+integer, pointer :: ptr(:)
+
+arr = [(i * 2, i=1,100)]
+
+!$omp target enter data map(to: arr(51:100))
+
+!$omp target
+arr(51:100) = arr(51:100) + 1
+!$omp end target
+
+!$omp target update from(arr(51:100:2))
+
+do i=1,100
+ if (i.le.50) then
+ if (arr(i).ne.i*2) stop 1
+ else
+ if (mod(i,2).eq.1 .and. arr(i).ne.i*2+1) stop 2
+ if (mod(i,2).eq.0 .and. arr(i).ne.i*2) stop 3
+ end if
+end do
+
+!$omp target exit data map(delete: arr)
+
+arr = [(i * 2, i=1,100)]
+
+! Similar, but update via pointer.
+
+ptr => arr(51:100)
+
+!$omp target enter data map(to: ptr(1:50))
+
+!$omp target
+ptr = ptr + 1
+!$omp end target
+
+!$omp target update from(ptr(::2))
+
+do i=1,100
+ if (i.le.50) then
+ if (arr(i).ne.i*2) stop 1
+ else
+ if (mod(i,2).eq.1 .and. arr(i).ne.i*2+1) stop 2
+ if (mod(i,2).eq.0 .and. arr(i).ne.i*2) stop 3
+ end if
+end do
+
+!$omp target exit data map(delete: ptr)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-6.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-6.f90
new file mode 100644
index 00000000000..5c42b9077b3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-6.f90
@@ -0,0 +1,34 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+program p
+implicit none
+integer, dimension(100) :: parr
+integer :: i
+
+parr = [(i, i=1,100)]
+
+!$omp target enter data map(to: parr)
+
+call s(parr)
+
+do i=1,100
+ if (mod(i,3).eq.1 .and. parr(i).ne.999) stop 1
+ if (mod(i,3).ne.1 .and. parr(i).ne.i) stop 2
+end do
+
+!$omp target exit data map(delete: parr)
+
+contains
+subroutine s(arr)
+implicit none
+integer, intent(inout) :: arr(*)
+
+!$omp target map(alloc: arr(1:100))
+arr(1:100) = 999
+!$omp end target
+
+!$omp target update from(arr(1:100:3))
+
+end subroutine s
+end program p
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-7.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-7.f90
new file mode 100644
index 00000000000..120fd9c90ed
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-7.f90
@@ -0,0 +1,36 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! Assumed-shape arrays
+
+program p
+implicit none
+integer, dimension(100) :: parr
+integer :: i
+
+parr = [(i, i=1,100)]
+
+!$omp target enter data map(to: parr)
+
+call s(parr)
+
+do i=1,100
+ if (mod(i,3).eq.1 .and. parr(i).ne.999) stop 1
+ if (mod(i,3).ne.1 .and. parr(i).ne.i) stop 2
+end do
+
+!$omp target exit data map(delete: parr)
+
+contains
+subroutine s(arr)
+implicit none
+integer, intent(inout) :: arr(:)
+
+!$omp target
+arr = 999
+!$omp end target
+
+!$omp target update from(arr(1:100:3))
+
+end subroutine s
+end program p
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-8.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-8.f90
new file mode 100644
index 00000000000..d9b3c9ca896
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-8.f90
@@ -0,0 +1,39 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! Test biasing for target-region lookup.
+
+implicit none
+integer, allocatable, target :: var(:,:,:)
+integer, pointer :: p(:,:,:)
+integer :: i, j, k
+
+allocate(var(1:20,5:25,10:30))
+
+var = 0
+
+!$omp target enter data map(to: var)
+
+!$omp target
+var = 99
+!$omp end target
+
+p => var(1:3:2,5:5,10:10)
+
+!$omp target update from(p)
+
+do i=1,20
+ do j=5,25
+ do k=10,30
+ if ((i.eq.1.or.i.eq.3).and.j.eq.5.and.k.eq.10) then
+ if (var(i,j,k).ne.99) stop 1
+ else
+ if (var(i,j,k).ne.0) stop 2
+ end if
+ end do
+ end do
+end do
+
+!$omp target exit data map(delete: var)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/noncontig-updates-9.f90
b/libgomp/testsuite/libgomp.fortran/noncontig-updates-9.f90
new file mode 100644
index 00000000000..689a46a91f0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/noncontig-updates-9.f90
@@ -0,0 +1,34 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! This test case hits the problem described in:
+! https://gcc.gnu.org/pipermail/gcc-patches/2023-February/612219.html
+
+! { dg-xfail-run-if "'enter data' bug" { offload_device_nonshared_as } }
+
+character(len=:), allocatable, dimension(:) :: lines
+integer :: i
+
+allocate(character(len=8) :: lines(10))
+
+lines = "OMPHELLO"
+
+!$omp target enter data map(to: lines)
+
+!$omp target
+lines = "NEWVALUE"
+!$omp end target
+
+!$omp target update from(lines(5:7:2))
+
+do i=1,10
+ if (i.eq.5.or.i.eq.7) then
+ if (lines(i).ne."NEWVALUE") stop 1
+ else
+ if (lines(i).ne."OMPHELLO") stop 2
+ end if
+end do
+
+!$omp target exit data map(delete: lines)
+
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target-update-iterators-4.f90
b/libgomp/testsuite/libgomp.fortran/target-update-iterators-4.f90
new file mode 100644
index 00000000000..a46e1f1c2e9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-update-iterators-4.f90
@@ -0,0 +1,36 @@
+! { dg-do run }
+! { dg-shouldfail "noncontiguous array section with iterator in target update"
}
+! { dg-output "Fortran runtime error: non-contiguous array section with
iterator in target update not supported" }
+
+! A genuinely noncontiguous (strided) array section combined with an
+! iterator on a "target update" cannot currently be lowered correctly
+! (iterator expansion conflicts with the noncontiguous/grid map lowering).
+! Check that this is rejected with a runtime error instead of silently
+! computing wrong results or crashing.
+
+program test
+ integer, parameter :: DIM1 = 4
+ integer, parameter :: DIM2 = 10
+
+ type :: array_ptr
+ integer, pointer :: arr(:)
+ end type
+
+ type (array_ptr) :: x(DIM1)
+ integer, target :: backing(DIM1, DIM2 * 2)
+ integer :: i, j
+
+ do i = 1, DIM1
+ do j = 1, DIM2 * 2
+ backing(i, j) = i * 100 + j
+ end do
+ x(i)%arr => backing(i, 1:DIM2 * 2:2)
+ end do
+
+ !$omp target enter data map(to: x)
+ !$omp target enter data map(to: backing)
+
+ !$omp target update to(iterator(i=1:DIM1): x(i)%arr(:))
+
+ print *, "unreachable"
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/target-update-iterators-5.f90
b/libgomp/testsuite/libgomp.fortran/target-update-iterators-5.f90
new file mode 100644
index 00000000000..d323ab735e2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-update-iterators-5.f90
@@ -0,0 +1,56 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+! { dg-xfail-run-if "GCC PR 127616: iterator dropped from target update" {
*-*-* } }
+
+! BUG producing wrong results with the following shape:
+!
+! target update to(iterator(i=1:N:100) : y(:, :, i, 2))
+!
+! Root cause: remove_unused_omp_iterator_vars (gimplify.cc) drops
+! OMP_CLAUSE_ITERATORS from a clause whenever it can't find a literal reference
+! to the iterator variable in the clause's DECL/SIZE trees, and only exempts
+! GOMP_MAP_GRID_DIM/GOMP_MAP_GRID_STRIDE clauses from that heuristic. For a
+! plain Fortran array, gfc_trans_omp_clauses evaluates the section's
+! address/length into scalar temporaries (gfc_evaluate_now) that no longer
+! contain a literal reference to the iterator variable -- and this update never
+! takes the grid path at all, so the exemption doesn't apply. The target
update
+! executes exactly once overall instead of once per iterator value, using a
+! bogus address.
+
+program test
+ implicit none
+ integer, parameter :: D1 = 4, D2 = 4, D3 = 9, D4 = 3
+ integer :: y(D1, D2, D3, D4)
+ integer :: i, j, k, errors
+
+ y = 0
+ !$omp target enter data map(to: y)
+
+ ! Set y(:,:,k,2) on the device only, for every k, to a value that
+ ! depends on k so a wrong (or missing) per-iteration update is
+ ! detectable.
+ !$omp target
+ do k = 1, D3
+ do j = 1, D2
+ do i = 1, D1
+ y(i, j, k, 2) = 100 * k + (j - 1) * D1 + i
+ end do
+ end do
+ end do
+ !$omp end target
+
+ ! Host copy is still all zero here. Pull each slab back individually
+ ! via the iterator, one k at a time.
+ !$omp target update from(iterator(k=1:D3): y(:, :, k, 2))
+
+ errors = 0
+ do k = 1, D3
+ do j = 1, D2
+ do i = 1, D1
+ if (y(i, j, k, 2) /= 100 * k + (j - 1) * D1 + i) errors = errors + 1
+ end do
+ end do
+ end do
+
+ if (errors /= 0) stop 1
+end program
--
2.55.0
