Hi Tobias,
See the attached fixup patches and my comments below.
On 14/09/2026 21:09, Tobias Burnus wrote:
Hi PA,
First, the following applies to Fortran and C/C++:
int a[10][10];
void f(int n) {
#pragma omp target update to(a[1][::-2])
}
This is accepted but violates:
* The stride must evaluate to a positive integer.
I thought we discussed this before and there was a check? Either I forgot
to apply some follow-up patch or I missed remembered?
The '[::<non-id-expr>' case was mishandled (see C++ sub-thread).
This now gives:
negative-stride.c:4:32: error: negative or zero stride in array section
in 'to' clause
4 | #pragma omp target update to(a[1][::-2])
In any case, For Fortran the same applies:
* If a stride expression is specified, it must be positive.
But there is no check either! See resolve_positive_int_expr for
how a check could look like (but it needs to be copied as the error
message is not quite right):
module m
implicit none
integer:: a(10, 10)
contains
subroutine sub (n)
integer, value :: n
integer :: i
!$omp target update to( a(::-2, 2))
end
end
Now handled as suggested:
10 | !$omp target update to( a(::-2, 2))
| 1
Error: Stride must be positive in TO clause at (1)
Added testcase bad-array-section-1.f90.
* * *
Paul-Antoine Arras wrote:
This patch implements noncontiguous "target update" for Fortran.
The existing middle end/runtime bits relating to C++ support are reused,
[...]
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.
There are three ways an array can be noncontiguous:
While going through testcases to ensure all three ways were properly
covered, I realised there was a measurable amount of C/C++ tests that
needed to be deduplicated or merged in addition to a few missing.
See below for which tests cover which of the three cases by language
(non-exhaustive list).
(A) There is an explicit stride in the array section - C/C++ and Fortran:
!$omp target update to( array2(::2) )
Here, the syntax is explicit, albeit for '(::N), it is only known
at runtime whether it is noncontiguous (N > 1) or not (N == 1,
assuming none of the below applies).
C: c-c++-common/gomp/array-section-3.c + libgomp.c/array-shaping-2.c
C++: g++.dg/gomp/array-section-7.C + libgomp.c++/array-shaping-6.C
Fortran: gfortran.dg/gomp/noncontig-updates-2.f90 +
libgomp.fortran/noncontig-updates-1.f90
(B) Only Fortran: The data object is intrinsically noncontiguous:
use iso_fortran_env
type t
integer(int128) :: x
integer :: y
end type
type(t) :: var(100)
integer(int128) :: array(100)
call sub(array(::2))
call sub(var(:)%x)
call sub(var(::2)%x)
...
subroutine sub(x)
integer(int128) :: x(:)
where 'X' is noncontiguous - and for the 'var' variants, as
c_sizeof(var(1))
is likely not a multiple of c_sizeof(var(1)%x), the stride offset is not
even
an integer number of the
such that already the plain
!$omp target update to( x ) / to(x(::)) / to(x(::1))
is internally noncontiguous.
...
gfortran.dg/gomp/noncontig-updates-5.f90 +
libgomp.fortran/noncontig-updates-4.f90
(C) Syntax wise contiguous, variable contiguous but still noncontiguous:
subroutine outer(N, M)
integer :: N, M
integer :: array(N,N)
integer, allocatable :: x(:,:)
allocate(x(N,N))
!$omp target update to(array(:M, :)) from(x(:M, : ))
or
void f(int N, int M, int *x)
{
int array[N][N];
!$omp target update to(array[:][:M])) from(([N][N])x[:][:M])
Here, 'x' and 'array' are contiguous – and the size is known. There is
also no stride specified explicitly, but for M < N (N > 1), the accessed
memory is noncontiguous.
C: gcc.dg/gomp/array-shaping-2.c + libgomp.c/array-shaping-2.c
C++: g++.dg/gomp/array-shaping-1.C + libgomp.c++/array-shaping-11.C
Fortran: gfortran.dg/gomp/noncontig-updates-3.f90 +
libgomp.fortran/noncontig-updates-11.f90
* * *
It seems as if all those cases need to work at runtime. However, I have not
re-checked the C cases (and not tried the Fortran cases at all).
* * *
Let's now add the iterator:
target update to(iterator(i=1:N) : array3[::2][i], array4[i][::2])
and, Fortran only:
target update to(iterator(i=1:N:100) : noncont[:][i:i+1])
and
target update to(iterator(i=1:N:100) : array([N][N])x[:M][i:i+1])
Except for the second case, it seems as if the strided-syntax support
plays no role in whether it is contiguous or not – and, hence, the
generated code needs to handle this for all cases – using the
non-contiguous code in the background.
There is no issue for
iterator(k=...) : var[2][k][:][:]
and
var[2][:n:m][:][:]
There *is* an issue. See below.
Here, everything after the '[k]' and '[:n:m]' is contiguous and either
the iterator or the strided code handles the discontinuity - assuming that
[:][:] are contiguous (they are in C but not in Fortran).
* * *
Thus, if an iterator is present:
* We can follow the C example and also in Fortran reject any explicitly
specified stride, unless compile-time known to be one
Yes, see below.
* Iterator followed by whole array sections is fine, i.e.
x[2][i][:][:] in C
y(:, :, i, 2) in Fortran
and can be accepted at compile time.
This fails both in Fortran and C -- but in different ways.
In Fortran, this is a pre-existing bug: remove_unused_omp_iterator_vars
drops the iterator because it only sees a scalar temporary without
direct reference to the iterator variable; see
libgomp/testsuite/libgomp.fortran/target-update-iterators-5.f90.
In C, this is specific to this patch: the per-iteration value of the
base address in GOMP_MAP_FROM_GRID/TO_GRID is not updated; see
libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-5.c.
Let me know if these two bugs should be fixed as part of this patch series.
* The following should be also OK in C/C++ or if simply contiguous in
Fortran:
x[i:], x[:i]
as every step is contiguous
This works:
libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-4.c.
* I believe the following is always mishandled, possibly except for
size=1 arrays:
x[n:m][i], y[:m][i], z[:][i]
and we can reject it.
Correct, this is checked by
gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c.
Thus, it seems as if only y(:, :, i, 2) needs a runtime check in Fortran.
For Fortran, I believe but have not tested it - that's all with 'iterator':
* Add a check for stride == 1 when explicitly specifying a stride and
otherwise always error (this needs to be added)
This...
* If gfc_is_not_contiguous determines that it is noncontiguous
error out (this needs to be added)
...and that are now implemented and checked in
gcc/testsuite/gfortran.dg/gomp/target-update-iterators-5.f90.
* Otherwise, use a run-time check (this is existing)
And for C, I think we need an error for:
x[n:m][i], y[:m][i], z[:][i]
(untested)
Confirmed: gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c.
and a runtime check for
x[2][i][:M][:]
for M == N, unless we always want to reject it?
As mentioned above, this always produces wrong results whatever the
value of M.
(BUILT_IN_GOMP_ERROR in gimplify.cc shows how one can produce
an error message.)
* * *
At least those are my initial thoughts.
The 'stride > 0' check is obvious. That there is no error shown in Fortran,
even if known to be noncontiguous need to be obviously fixed as well.
For iterator and any '::n' access (stride != 1 (intcst)), I guess we want
to follow C/C++ and also error out unconditionally.
And for the rest: I have not checked, but it seems as if some cases should
give either a compile- or runtime error in C/C++ and Fortran, but I think
Fortran's current runtime check handles all those, such that only something
for C/C++ needs to be done (untested).
* * *
openmp.cc: besides the stride >= 1 check, there is a tailing ' ' in
the line '&& code->op != EXEC_OACC_UPDATE'
Fixed.
* * *
Otherwise for completeness:
* expr.cc is fine
* gfc_omp_calculate_gcd is OK, gfc_trans_omp_arrayshape_type seems to be
fine, but I still have to look at its use.
TODO on my side:
I also still need to digest the rest of trans-openmp.cc and in how far
that handles/should handle my
type t
integer(int128) :: x
integer :: y
end type
type(t) :: var(100)
call sub(var(::2)%x)
where 'type' likely has size 16 + 4 = 20 bytes, i.e stride multiplier = 40
byte – while 'x' being 16 bytes can only represent the strides 16, 32,
48, 64, ...
And, if it not handles it, how should it handle it. (There is an issue
with the
internal representation in gfortran that eventually needs to be fixed -
but we shouldn't
repeat the mistake here.
The proper way is lower bound, extent and stride multiplier - gfortran
uses lower bound,
upper bound [= max(0, upper - lower + 1)], stride, offset, and as band-
aid 'span', at least
sometimes.)
Tobias
Thanks,
--
PA
From 9e373453753bfb72738a636a50b13ac9046ecb46 Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <[email protected]>
Date: Wed, 23 Sep 2026 18:47:55 +0200
Subject: [PATCH 6/6] ME: relax grid-iterators check
---
gcc/gimplify.cc | 16 ++-
.../gomp/target-update-iterators-4.c | 27 ++++-
.../gomp/target-update-iterators-6.c | 34 ++++++
.../target-update-iterators-4.c | 90 ++++++++++++++
.../target-update-iterators-5.c | 112 ++++++++++++++++++
5 files changed, 275 insertions(+), 4 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterators-6.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-4.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-5.c
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 9dfa54fb69e..62dea6dad2a 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -19191,11 +19191,23 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE))
{
- sorry_at (OMP_CLAUSE_LOCATION (c),
+ for (tree it = OMP_CLAUSE_ITERATORS (c); it; it = TREE_CHAIN (it))
+ {
+ tree var = OMP_ITERATOR_VAR (it);
+ tree t = walk_tree (&OMP_CLAUSE_DECL (c), find_var_decl, var, NULL);
+ if (t == NULL_TREE)
+ t = walk_tree (&OMP_CLAUSE_SIZE (c), find_var_decl, var, NULL);
+ if (t != NULL_TREE)
+ {
+ sorry_at (
+ OMP_CLAUSE_LOCATION (c),
"%<iterator%> modifier combined with a non-contiguous "
"array section");
- break;
+ goto done;
+ }
+ }
}
+done:
build_omp_iterators_loops (&OMP_STANDALONE_CLAUSES (expr),
&iterator_loops_seq);
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c
index 8f18d6b0ff7..8a2630eba02 100644
--- a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-4.c
@@ -1,8 +1,7 @@
/* { dg-do compile } */
/* { dg-options "-fopenmp" } */
-/* Check that a map iterator on a strided target update is rejected with a sorry
- message. */
+/* Combinations of map iterators and array sections that are rejected. */
#define DIM1 17
#define DIM2 39
@@ -11,3 +10,27 @@ void f (int **x)
{
#pragma omp target update to (iterator(i=0:DIM1): x[i][0:DIM2:2]) /* { dg-message "sorry, unimplemented: 'iterator' modifier combined with a non-contiguous array section" } */
}
+
+void f1 (int m)
+{
+ int x[10][10];
+ #pragma omp target update to(iterator(i=0:9): x[0:m][i]) /* { dg-message "sorry, unimplemented: 'iterator' modifier combined with a non-contiguous array section" } */
+}
+
+void f2 (int m)
+{
+ int y[10][10];
+ #pragma omp target update to(iterator(i=0:9): y[:m][i]) /* { dg-message "sorry, unimplemented: 'iterator' modifier combined with a non-contiguous array section" } */
+}
+
+void f3 (void)
+{
+ int z[10][10];
+ #pragma omp target update to(iterator(i=0:9): z[:][i]) /* { dg-message "sorry, unimplemented: 'iterator' modifier combined with a non-contiguous array section" } */
+}
+
+void f4 (void)
+{
+ int x[10][10];
+ #pragma omp target update to(iterator(i=0:9): x[::2][i]) /* { dg-message "sorry, unimplemented: 'iterator' modifier combined with a non-contiguous array section" } */
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-6.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-6.c
new file mode 100644
index 00000000000..3014ec5d085
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-6.c
@@ -0,0 +1,34 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+/* Check allowed and rejected combinations of map iterators and array
+ sections. */
+
+void f1 (void)
+{
+ int x[3][10][10][10];
+ #pragma omp target update to(iterator(i=0:9): x[2][i][:][:]) /* { dg-message "sorry, unimplemented: 'iterator' modifier combined with a non-contiguous array section" } */
+}
+
+void f2 (int m)
+{
+ int x[3][10][10][10];
+ #pragma omp target update to(iterator(i=0:9): x[2][i][:m][:]) /* { dg-message "sorry, unimplemented: 'iterator' modifier combined with a non-contiguous array section" } */
+}
+
+
+void f3 (void)
+{
+ int x[10];
+ #pragma omp target update to(iterator(i=0:9): x[i:])
+}
+
+void f4 (void)
+{
+ int x[10];
+ #pragma omp target update to(iterator(i=0:9): x[:i])
+}
+
+/* f3/f4 (iterator used as a section bound, not a separate index) never take
+ the grid path at all -- checked directly below. */
+/* { dg-final { scan-tree-dump-times "update to\\(iterator\\(int i=0:9:1" 2 "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-4.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-4.c
new file mode 100644
index 00000000000..6f8d61d271d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-4.c
@@ -0,0 +1,90 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* The iterator variable is used as a section bound (x[i:], x[:i]), rather than
+ as a separate index. Each iteration is a single contiguous run, so no
+ grid/noncontiguous machinery is involved; check that every iteration resolves
+ the bound using its own value of the iterator. Choosing iterator ranges that
+ stop short of the array's other end (instead of spanning the whole array)
+ makes that checkable: positions outside every iteration's range must stay
+ untouched. */
+
+#define DIM1 10
+
+int x[DIM1];
+
+void
+test_suffix_bound (void)
+{
+ int errors;
+
+ for (int j = 0; j < DIM1; j++)
+ x[j] = 1000 + j;
+
+#pragma omp target enter data map(to: x)
+
+#pragma omp target
+ for (int j = 0; j < DIM1; j++)
+ x[j] = 9000 + j;
+
+ for (int j = 0; j < DIM1; j++)
+ x[j] = -1;
+
+ /* Largest section copied is at i=5 (x[5:], positions 5..9); positions
+ 0..4 must not be touched. */
+#pragma omp target update from(iterator(i=5:10): x[i:])
+
+ errors = 0;
+ for (int j = 0; j < 5; j++)
+ if (x[j] != -1)
+ errors++;
+ for (int j = 5; j < DIM1; j++)
+ if (x[j] != 9000 + j)
+ errors++;
+
+ if (errors != 0)
+ __builtin_abort ();
+}
+
+int y[DIM1];
+
+void
+test_prefix_bound (void)
+{
+ int errors;
+
+ for (int j = 0; j < DIM1; j++)
+ y[j] = 2000 + j;
+
+#pragma omp target enter data map(to: y)
+
+#pragma omp target
+ for (int j = 0; j < DIM1; j++)
+ y[j] = 8000 + j;
+
+ for (int j = 0; j < DIM1; j++)
+ y[j] = -1;
+
+ /* x[:i] has length i, so the largest section copied is at i=5 (y[:5],
+ positions 0..4); positions 5..9 must not be touched. */
+#pragma omp target update from(iterator(i=1:6): y[:i])
+
+ errors = 0;
+ for (int j = 0; j < 5; j++)
+ if (y[j] != 8000 + j)
+ errors++;
+ for (int j = 5; j < DIM1; j++)
+ if (y[j] != -1)
+ errors++;
+
+ if (errors != 0)
+ __builtin_abort ();
+}
+
+int
+main (void)
+{
+ test_suffix_bound ();
+ test_prefix_bound ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-5.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-5.c
new file mode 100644
index 00000000000..35c86ac03a4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-5.c
@@ -0,0 +1,112 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* BUG: both x[2][i][:][:] and x[2][i][:m][:] are mishandled and produce wrong
+ results.
+
+ Root cause: the expanded iterator loop only stores a per-iteration value into
+ the elems array for the GOMP_MAP_GRID_DIM clauses' iterator; the
+ GOMP_MAP_FROM_GRID/TO_GRID clause's own elems slot is written once, before
+ the loop, to a bogus constant, and never updated inside the loop body. So
+ the per-iteration base address is never actually captured per iteration.
+ Confirmed via a control case: the identical shape with a fixed index instead
+ of an iterator (x[2][5][:][:], no iterator at all) updates correctly. */
+
+#define DIM 10
+
+int x[3][DIM][DIM][DIM];
+
+static void
+reset_and_upload (int fill)
+{
+ for (int p = 0; p < 3; p++)
+ for (int i = 0; i < DIM; i++)
+ for (int j = 0; j < DIM; j++)
+ for (int k = 0; k < DIM; k++)
+ x[p][i][j][k] = fill;
+
+#pragma omp target enter data map(to: x)
+}
+
+static void
+poison_host (void)
+{
+ for (int p = 0; p < 3; p++)
+ for (int i = 0; i < DIM; i++)
+ for (int j = 0; j < DIM; j++)
+ for (int k = 0; k < DIM; k++)
+ x[p][i][j][k] = -1;
+}
+
+void
+test_whole_trailing_sections (void)
+{
+ int errors;
+
+ reset_and_upload (0);
+
+#pragma omp target
+ for (int i = 0; i < DIM; i++)
+ for (int j = 0; j < DIM; j++)
+ for (int k = 0; k < DIM; k++)
+ x[2][i][j][k] = 9000 + 100 * i + 10 * j + k;
+
+ poison_host ();
+
+#pragma omp target update from(iterator(i=0:DIM): x[2][i][:][:])
+
+ errors = 0;
+ for (int p = 0; p < 3; p++)
+ for (int i = 0; i < DIM; i++)
+ for (int j = 0; j < DIM; j++)
+ for (int k = 0; k < DIM; k++)
+ {
+ int expected = (p == 2) ? 9000 + 100 * i + 10 * j + k : -1;
+ if (x[p][i][j][k] != expected)
+ errors++;
+ }
+
+ if (errors != 0)
+ __builtin_abort ();
+}
+
+void
+test_partial_trailing_section (int m)
+{
+ int errors;
+
+ reset_and_upload (0);
+
+#pragma omp target
+ for (int i = 0; i < DIM; i++)
+ for (int j = 0; j < DIM; j++)
+ for (int k = 0; k < DIM; k++)
+ x[2][i][j][k] = 8000 + 100 * i + 10 * j + k;
+
+ poison_host ();
+
+#pragma omp target update from(iterator(i=0:DIM): x[2][i][:m][:])
+
+ errors = 0;
+ for (int p = 0; p < 3; p++)
+ for (int i = 0; i < DIM; i++)
+ for (int j = 0; j < DIM; j++)
+ for (int k = 0; k < DIM; k++)
+ {
+ int expected = (p == 2 && j < m)
+ ? 8000 + 100 * i + 10 * j + k : -1;
+ if (x[p][i][j][k] != expected)
+ errors++;
+ }
+
+ if (errors != 0)
+ __builtin_abort ();
+}
+
+int
+main (void)
+{
+ test_whole_trailing_sections ();
+ test_partial_trailing_section (4);
+ return 0;
+}
--
2.55.0
From 2950e9bc0b007a8ef0d75bfd4d98cf555390d502 Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <[email protected]>
Date: Wed, 23 Sep 2026 18:44:44 +0200
Subject: [PATCH 5/6] Fortran: target update with map iterators
---
gcc/fortran/openmp.cc | 17 ++++--
.../gomp/target-update-iterators-5.f90 | 23 ++++++++
.../gomp/target-update-iterators-6.f90 | 39 +++++++++++++
.../target-update-iterators-5.f90 | 55 +++++++++++++++++++
4 files changed, 130 insertions(+), 4 deletions(-)
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/target-update-iterators-5.f90
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index bf687d985bd..177b0a4ecae 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -10209,6 +10209,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
@@ -10260,14 +10266,17 @@ resolve_omp_clauses_aff_dep_map_cache (gfc_code *code,
"in %s clause at %L", name, &n->where);
break;
}
- else if (ar->stride[i] && code
+ 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 (ar->stride[i]->expr_type == EXPR_CONSTANT
- && ar->stride[i]->ts.type == BT_INTEGER
- && mpz_sgn (ar->stride[i]->value.integer) <= 0)
+ 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);
+ 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
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/target-update-iterators-5.f90 b/libgomp/testsuite/libgomp.fortran/target-update-iterators-5.f90
new file mode 100644
index 00000000000..abd37937f46
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-update-iterators-5.f90
@@ -0,0 +1,55 @@
+! { dg-do run }
+! { dg-require-effective-target offload_device_nonshared_as }
+
+! 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
From 30ace11f2bdb7820f4742c584cac9fd75b26fa66 Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <[email protected]>
Date: Tue, 22 Sep 2026 16:18:30 +0200
Subject: [PATCH 4/6] ME: fix handling of fixed indices mixed with array
sections
---
gcc/omp-general.cc | 13 ++-
.../c-c++-common/gomp/target-update-8.c | 26 ++++++
.../target-update-leading-index-1.c | 81 +++++++++++++++++++
3 files changed, 116 insertions(+), 4 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-8.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-update-leading-index-1.c
diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc
index f22143881b8..4dca2bd35cb 100644
--- a/gcc/omp-general.cc
+++ b/gcc/omp-general.cc
@@ -4207,6 +4207,7 @@ static bool
omp_parse_noncontiguous_array (tree *expr0)
{
tree expr = *expr0;
+ tree noncontig_base = expr;
bool noncontig = false;
while (TREE_CODE (expr) == OMP_ARRAY_SECTION
@@ -4215,14 +4216,18 @@ omp_parse_noncontiguous_array (tree *expr0)
/* Contiguous arrays use ARRAY_REF. By the time we reach here,
OMP_ARRAY_SECTION is only used for noncontiguous arrays. */
if (TREE_CODE (expr) == OMP_ARRAY_SECTION)
- noncontig = true;
-
- expr = TREE_OPERAND (expr, 0);
+ {
+ noncontig = true;
+ expr = TREE_OPERAND (expr, 0);
+ noncontig_base = expr;
+ }
+ else
+ expr = TREE_OPERAND (expr, 0);
}
if (noncontig)
{
- *expr0 = expr;
+ *expr0 = noncontig_base;
return true;
}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-8.c b/gcc/testsuite/c-c++-common/gomp/target-update-8.c
new file mode 100644
index 00000000000..85398d5b786
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-8.c
@@ -0,0 +1,26 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+/* Check that fixed indices mixed with array sections are lowered correctly. */
+
+int x[3][10];
+
+void
+leading_index (void)
+{
+#pragma omp target update from(x[1][0:5:2])
+}
+
+
+int y[20][20][20];
+
+void
+interspersed_index (void)
+{
+#pragma omp target update from(y[0:2:2][1][0:2:2])
+}
+
+/* { dg-final { scan-tree-dump "map\\(from_grid:x\\\[1\\\] \\\[len: 4\\\]\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "map\\(from_grid:y \\\[len: 4\\\]\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(grid_dim:" 4 "gimple" } } */
+/* { dg-final { scan-tree-dump "map\\(grid_dim:1 \\\[len: 1\\\]\\)" "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-leading-index-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-leading-index-1.c
new file mode 100644
index 00000000000..db0e8c668a9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-leading-index-1.c
@@ -0,0 +1,81 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* Check that fixed indices mixed with array sections work correctly. */
+
+int x[3][10];
+
+void
+test_leading_index (void)
+{
+ int j, errors;
+
+ for (j = 0; j < 10; j++)
+ {
+ x[0][j] = 1000 + j;
+ x[1][j] = 2000 + j;
+ x[2][j] = 3000 + j;
+ }
+
+#pragma omp target enter data map(to: x)
+
+#pragma omp target
+ for (j = 0; j < 10; j += 2)
+ x[1][j] = 9000 + j;
+
+ for (j = 0; j < 30; j++)
+ ((int *) x)[j] = -1;
+
+#pragma omp target update from(x[1][0:5:2])
+
+ errors = 0;
+ for (j = 0; j < 10; j += 2)
+ if (x[1][j] != 9000 + j)
+ errors++;
+
+ if (errors != 0)
+ __builtin_abort ();
+}
+
+
+int y[20][20][20];
+
+void
+test_interspersed_index (void)
+{
+ int i, k, errors;
+
+ for (i = 0; i < 20; i++)
+ for (int j = 0; j < 20; j++)
+ for (k = 0; k < 20; k++)
+ y[i][j][k] = i * 1000 + j * 100 + k;
+
+#pragma omp target enter data map(to: y)
+
+#pragma omp target
+ for (i = 0; i < 4; i += 2)
+ for (k = 0; k < 4; k += 2)
+ y[i][1][k] = 9000 + i * 10 + k;
+
+ for (i = 0; i < 8000; i++)
+ ((int *) y)[i] = -1;
+
+#pragma omp target update from(y[0:2:2][1][0:2:2])
+
+ errors = 0;
+ for (i = 0; i < 4; i += 2)
+ for (k = 0; k < 4; k += 2)
+ if (y[i][1][k] != 9000 + i * 10 + k)
+ errors++;
+
+ if (errors != 0)
+ __builtin_abort ();
+}
+
+int
+main (void)
+{
+ test_leading_index ();
+ test_interspersed_index ();
+ return 0;
+}
--
2.55.0
From 083d184e8898a018de4cbe7a6eb4c579b3da9326 Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <[email protected]>
Date: Mon, 21 Sep 2026 15:58:46 +0200
Subject: [PATCH 3/6] Fortran: missing testcases
---
.../gfortran.dg/gomp/bad-array-section-1.f90 | 31 +++++++++++++++++++
.../gfortran.dg/gomp/noncontig-updates-5.f90 | 28 +++++++++++++++++
.../gfortran.dg/gomp/noncontig-updates-6.f90 | 15 +++++++++
3 files changed, 74 insertions(+)
create mode 100644 gcc/testsuite/gfortran.dg/gomp/bad-array-section-1.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
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-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
--
2.55.0
From 077ecf3f5ddd25dfb257dd6c54c1ec3f345c7876 Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <[email protected]>
Date: Mon, 21 Sep 2026 15:44:52 +0200
Subject: [PATCH 2/6] C/C++: Dedup and merge + add missing testcases
---
.../c-c++-common/gomp/array-section-3.c | 18 ++
.../c-c++-common/gomp/array-shaping-1.c | 14 ++
.../gomp/bad-array-shaping-2.c} | 9 +-
.../gomp/bad-array-shaping-3.c} | 7 +-
gcc/testsuite/g++.dg/gomp/array-section-7.C | 9 -
.../g++.dg/gomp/bad-array-shaping-1.C | 14 --
.../g++.dg/gomp/bad-array-shaping-8.C | 20 --
gcc/testsuite/gcc.dg/gomp/array-shaping-2.c | 16 ++
.../testsuite/libgomp.c++/array-shaping-1.C | 226 ------------------
.../testsuite/libgomp.c++/array-shaping-10.C | 22 --
.../testsuite/libgomp.c++/array-shaping-12.C | 24 --
.../testsuite/libgomp.c++/array-shaping-13.C | 35 ---
.../testsuite/libgomp.c++/array-shaping-5.C | 38 ---
.../libgomp.c-c++-common/array-shaping-15.c | 41 ++++
.../array-shaping-17.c} | 18 +-
.../array-shaping-18.c} | 2 +-
.../array-shaping-19.c} | 0
.../array-shaping-20.c} | 0
.../array-shaping-21.c} | 2 +-
19 files changed, 109 insertions(+), 406 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-section-3.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/array-shaping-1.c
rename gcc/testsuite/{gcc.dg/gomp/bad-array-shaping-c-1.c => c-c++-common/gomp/bad-array-shaping-2.c} (55%)
rename gcc/testsuite/{gcc.dg/gomp/bad-array-shaping-c-6.c => c-c++-common/gomp/bad-array-shaping-3.c} (83%)
create mode 100644 gcc/testsuite/gcc.dg/gomp/array-shaping-2.c
delete mode 100644 libgomp/testsuite/libgomp.c++/array-shaping-5.C
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/array-shaping-15.c
rename libgomp/testsuite/{libgomp.c++/array-shaping-2.C => libgomp.c-c++-common/array-shaping-17.c} (68%)
rename libgomp/testsuite/{libgomp.c/array-shaping-1.c => libgomp.c-c++-common/array-shaping-18.c} (99%)
rename libgomp/testsuite/{libgomp.c/array-shaping-4.c => libgomp.c-c++-common/array-shaping-19.c} (100%)
rename libgomp/testsuite/{libgomp.c/array-shaping-5.c => libgomp.c-c++-common/array-shaping-20.c} (100%)
rename libgomp/testsuite/{libgomp.c/array-shaping-6.c => libgomp.c-c++-common/array-shaping-21.c} (94%)
diff --git a/gcc/testsuite/c-c++-common/gomp/array-section-3.c b/gcc/testsuite/c-c++-common/gomp/array-section-3.c
new file mode 100644
index 00000000000..ff2698cb1de
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/array-section-3.c
@@ -0,0 +1,18 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+/* An explicit stride makes the section noncontiguous. Check that both a literal
+ and a variable stride route through the noncontiguous-update lowering:
+ to_grid, grid_dim and grid_stride. */
+
+int arr[128];
+
+void
+g (int s)
+{
+#pragma omp target update to(arr[::2])
+ /* { dg-final { scan-tree-dump {map\(to_grid:arr \[len: 4\]\) map\(grid_dim:0 \[len: 64\]\) map\(grid_stride:2\)} "original" } } */
+
+#pragma omp target update to(arr[::(s)])
+ /* { dg-final { scan-tree-dump {map\(to_grid:arr \[len: 4\]\) map\(grid_dim:0 \[len: 128 /\[cl\] \(sizetype\) s\]\) map\(grid_stride:s\)} "original" } } */
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/array-shaping-1.c b/gcc/testsuite/c-c++-common/gomp/array-shaping-1.c
new file mode 100644
index 00000000000..e8181b1b044
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/array-shaping-1.c
@@ -0,0 +1,14 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+/* The reshaped array is contiguous overall and the section has no explicit
+ stride, but for runtime-variable bounds "f"/"g" the accessed memory may be
+ noncontiguous.
+ Check that this routes through the grid/noncontiguous-update lowering. */
+
+void
+foo (float *w, int e, int f, int g)
+{
+#pragma omp target update to(([e][10]) w[3:2][f:g])
+ /* { dg-final { scan-tree-dump {map\(to_grid:VIEW_CONVERT_EXPR.*\(\*w\) \[len: 4\]\) map\(grid_dim:3 \[len: 2\]\) map\(grid_dim:[^\[]+\[len: g\]\)} "original" } } */
+}
diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c b/gcc/testsuite/c-c++-common/gomp/bad-array-shaping-2.c
similarity index 55%
rename from gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c
rename to gcc/testsuite/c-c++-common/gomp/bad-array-shaping-2.c
index 42d584fa624..954545cf9d4 100644
--- a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/bad-array-shaping-2.c
@@ -1,12 +1,11 @@
-// { dg-do compile }
+/* { dg-do compile } */
#include <string.h>
-#include <assert.h>
#include <stdlib.h>
int main (void)
{
- float *arr = calloc (100, sizeof (float));
+ float *arr = (float *) calloc (100, sizeof (float));
#pragma omp target enter data map(to: arr[:100])
@@ -15,8 +14,8 @@ int main (void)
arr[j * 10 + i] = i + j * 3;
#pragma omp target update to(([10][10]) arr[3:2][1:8][0:5])
-// { dg-error "too many array section specifiers for" "" { target *-*-* } .-1 }
-// { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 }
+/* { dg-error "too many array section specifiers for" "" { target *-*-* } .-1 } */
+/* { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } */
#pragma omp target exit data map(from: arr[:100])
diff --git a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-6.c b/gcc/testsuite/c-c++-common/gomp/bad-array-shaping-3.c
similarity index 83%
rename from gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-6.c
rename to gcc/testsuite/c-c++-common/gomp/bad-array-shaping-3.c
index 7de4323ee35..bfcd85ae628 100644
--- a/gcc/testsuite/gcc.dg/gomp/bad-array-shaping-c-6.c
+++ b/gcc/testsuite/c-c++-common/gomp/bad-array-shaping-3.c
@@ -1,12 +1,13 @@
-// { dg-do compile }
+/* { dg-do compile } */
int main (void)
{
char *ptr;
#pragma omp target update to(([5][6][7]) ptr[0:4][0:7][0:7])
-/* { dg-error {length '7' above array section size in 'to' clause} "" { target *-*-* } .-1 } */
-/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } */
+/* { dg-error {length '7' above array section size in 'to' clause} "" { target c } .-1 } */
+/* { dg-error {length '7' with stride '1' above array section size in 'to' clause} "" { target c++ } .-2 } */
+/* { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-3 } */
#pragma omp target update to(([5][6][7]) ptr[1:5][0:6][0:7])
/* { dg-error {high bound '6' above array section size in 'to' clause} "" { target *-*-* } .-1 } */
diff --git a/gcc/testsuite/g++.dg/gomp/array-section-7.C b/gcc/testsuite/g++.dg/gomp/array-section-7.C
index 9fb282bfde4..70973136b10 100644
--- a/gcc/testsuite/g++.dg/gomp/array-section-7.C
+++ b/gcc/testsuite/g++.dg/gomp/array-section-7.C
@@ -32,15 +32,6 @@ void g() {
#pragma omp target update to(arr[::])
// { dg-final { scan-tree-dump "to\\(arr\\\[0\\\] \\\[len: 512\\\]\\)" "original" } }
- // Valid: no ambiguity, since a numeric literal can never start an
- // id-expression, so "2" is unambiguously the stride.
-#pragma omp target update to(arr[::2])
-// { dg-final { scan-tree-dump "map\\(to_grid:arr \\\[len: 4\\\]\\) map\\(grid_dim:0 \\\[len: 64\\\]\\) map\\(grid_stride:2\\)" "original" } }
-
- // Likewise for a parenthesized expression as the stride.
-#pragma omp target update to(arr[::(s)])
-// { dg-final { scan-tree-dump "map\\(to_grid:arr \\\[len: 4\\\]\\) map\\(grid_dim:0 \\\[len: 128 /\\\[cl\\\] \\(sizetype\\) s\\\]\\) map\\(grid_stride:s\\)" "original" } }
-
// A negative or zero stride is rejected as a semantic error (not a parse
// failure), same as when written with a disambiguating space.
#pragma omp target update to(arr[::-2])
diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-1.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-1.C
index 84bd729a939..ecb1f44c3c7 100644
--- a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-1.C
+++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-1.C
@@ -25,20 +25,6 @@ int main()
{
float *arr = new float[100];
- memset (arr, 0, sizeof (float) * 100);
-
-#pragma omp target enter data map(to: arr[ :100])
-
- for (int j = 0; j < 10; j++)
- for (int i = 0; i < 10; i++)
- arr[j * 10 + i] = i + j * 3;
-
-#pragma omp target update to(([10][10]) arr[3:2][1:8][0:5])
-// { dg-error "too many array section specifiers for" "" { target *-*-* } .-1 }
-// { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 }
-
-#pragma omp target exit data map(from: arr[ :100])
-
foo<float, 5, 20> (arr);
delete[] arr;
diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C
index 4b077ca86a9..e61a06091c3 100644
--- a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C
+++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C
@@ -24,26 +24,6 @@ void foo ()
int main()
{
- char *ptr;
-
-#pragma omp target update to(([5][6][7]) ptr[0:4][0:7][0:7])
-// { dg-error {length '7' with stride '1' above array section size in 'to' clause} "" { target *-*-* } .-1 }
-// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 }
-
-#pragma omp target update to(([5][6][7]) ptr[1:5][0:6][0:7])
-// { dg-error {high bound '6' above array section size in 'to' clause} "" { target *-*-* } .-1 }
-// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 }
-
-#pragma omp target update from(([100]) ptr[3:33:3])
-
-#pragma omp target update from(([100]) ptr[4:33:3])
-// { dg-error {high bound '101' above array section size in 'from' clause} "" { target *-*-* } .-1 }
-// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 }
-
-#pragma omp target update to(([10][10]) ptr[0:9:-1][0:9])
-// { dg-error {negative or zero stride in array section in 'to' clause} "" { target *-*-* } .-1 }
-// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 }
-
foo<char> ();
return 0;
diff --git a/gcc/testsuite/gcc.dg/gomp/array-shaping-2.c b/gcc/testsuite/gcc.dg/gomp/array-shaping-2.c
new file mode 100644
index 00000000000..67d5f1b2fc2
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/array-shaping-2.c
@@ -0,0 +1,16 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+/* The reshaped array is contiguous overall and the section has no explicit
+ stride, but for a runtime-variable "len" the accessed memory may be
+ noncontiguous.
+ Check that this routes through the grid/noncontiguous-update lowering.
+ Check in particular that VIEW_CONVERT_EXPR applies to the bare pointer "a",
+ which is specific to the C front end. */
+
+void
+array_shaping (double *a, int nx, int ny, int lo, int len)
+{
+#pragma omp target update from((([nx][ny]) a)[0:nx][lo:len])
+ /* { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR.*\(a\) \[len: 8\]\) map\(grid_dim:0 \[len: nx\]\) map\(grid_dim:[^\[]+\[len: len\]\)} "original" } } */
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-1.C b/libgomp/testsuite/libgomp.c++/array-shaping-1.C
index 6ff5f9475f6..5e0857464c7 100644
--- a/libgomp/testsuite/libgomp.c++/array-shaping-1.C
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-1.C
@@ -237,232 +237,6 @@ void foo()
int main()
{
- int *arr;
- int x = 5;
- int arr2d[10][10];
-
- arr = new int[100];
-
- /* Update whole reshaped array. */
-
- memset (arr, 0, 100 * sizeof (int));
-
-#pragma omp target enter data map(to: arr[:100])
-
- for (int j = 0; j < x; j++)
- for (int i = 0; i < 10; i++)
- arr[j * 10 + i] = i ^ j;
-
-#pragma omp target update to(([10][x]) arr)
-
-#pragma omp target exit data map(from: arr[:100])
-
- for (int j = 0; j < 10; j++)
- for (int i = 0; i < 10; i++)
- if (j < x)
- assert (arr[j * 10 + i] == i ^ j);
- else
- assert (arr[j * 10 + i] == 0);
-
-
- /* Strided update. */
-
- memset (arr, 0, 100 * sizeof (int));
-
-#pragma omp target enter data map(to: arr[:100])
-
- for (int j = 0; j < 20; j++)
- for (int i = 0; i < 5; i++)
- arr[j * 5 + i] = i + j;
-
-#pragma omp target update to(([5][5]) arr[0:3][0:3:2])
-
-#pragma omp target exit data map(from: arr[:100])
-
- for (int j = 0; j < 20; j++)
- for (int i = 0; i < 5; i++)
- if (j < 3 && (i & 1) == 0 && i < 6)
- assert (arr[j * 5 + i] == i + j);
- else
- assert (arr[j * 5 + i] == 0);
-
-
- /* Reshaped update, contiguous. */
-
- memset (arr, 0, 100 * sizeof (int));
-
-#pragma omp target enter data map(to: arr[:100])
-
- for (int j = 0; j < 20; j++)
- for (int i = 0; i < 5; i++)
- arr[j * 5 + i] = 2 * j + i;
-
-#pragma omp target update to(([5][5]) arr[0:5][0:5])
-
-#pragma omp target exit data map(from: arr[:100])
-
- for (int j = 0; j < 20; j++)
- for (int i = 0; i < 5; i++)
- if (j < 5 && i < 5)
- assert (arr[j * 5 + i] == 2 * j + i);
- else
- assert (arr[j * 5 + i] == 0);
-
-
- /* Strided update on actual array. */
-
- memset (arr2d, 0, 100 * sizeof (int));
-
-#pragma omp target enter data map(to: arr2d)
-
- for (int j = 0; j < 10; j++)
- for (int i = 0; i < 10; i++)
- arr2d[j][i] = j + 2 * i;
-
-#pragma omp target update to(arr2d[0:5:2][5:2])
-
-#pragma omp target exit data map(from: arr2d)
-
- for (int j = 0; j < 10; j++)
- for (int i = 0; i < 10; i++)
- if ((j & 1) == 0 && i >= 5 && i < 7)
- assert (arr2d[j][i] == j + 2 * i);
- else
- assert (arr2d[j][i] == 0);
-
-
- /* Update with non-constant bounds. */
-
- memset (arr, 0, 100 * sizeof (int));
-
-#pragma omp target enter data map(to: arr[:100])
-
- for (int j = 0; j < 10; j++)
- for (int i = 0; i < 10; i++)
- arr[j * 10 + i] = (2 * j) ^ i;
-
- x = 3;
- int y = yy, z = zz, str = str_str;
- /* This is actually [0:3:2] [4:2:2]. */
-#pragma omp target update to(([10][10]) arr[0:x:2][y:z:str])
-
-#pragma omp target exit data map(from: arr[:100])
-
- for (int j = 0; j < 10; j++)
- for (int i = 0; i < 10; i++)
- if ((j & 1) == 0 && j < 6 && (i & 1) == 0 && i >= 4 && i < 8)
- assert (arr[j * 10 + i] == (2 * j) ^ i);
- else
- assert (arr[j * 10 + i] == 0);
-
-
- /* Update with full "major" dimension. */
-
- memset (arr, 0, 100 * sizeof (int));
-
-#pragma omp target enter data map(to: arr[:100])
-
- for (int j = 0; j < 10; j++)
- for (int i = 0; i < 10; i++)
- arr[j * 10 + i] = i + j;
-
-#pragma omp target update to(([10][10]) arr[0:10][3:1])
-
-#pragma omp target exit data map(from: arr[:100])
-
- for (int j = 0; j < 10; j++)
- for (int i = 0; i < 10; i++)
- if (i == 3)
- assert (arr[j * 10 + i] == i + j);
- else
- assert (arr[j * 10 + i] == 0);
-
-
- /* Update with full "minor" dimension. */
-
- memset (arr, 0, 100 * sizeof (int));
-
-#pragma omp target enter data map(to: arr[:100])
-
- for (int j = 0; j < 10; j++)
- for (int i = 0; i < 10; i++)
- arr[j * 10 + i] = 3 * (i + j);
-
-#pragma omp target update to(([10][10]) arr[3:2][0:10])
-
-#pragma omp target exit data map(from: arr[:100])
-
- for (int j = 0; j < 10; j++)
- for (int i = 0; i < 10; i++)
- if (j >= 3 && j < 5)
- assert (arr[j * 10 + i] == 3 * (i + j));
- else
- assert (arr[j * 10 + i] == 0);
-
-
- /* Rectangle update. */
-
- memset (arr, 0, 100 * sizeof (int));
-
-#pragma omp target enter data map(to: arr[:100])
-
- for (int j = 0; j < 10; j++)
- for (int i = 0; i < 10; i++)
- arr[j * 10 + i] = 5 * (i + j);
-
-#pragma omp target update to(([10][10]) arr[3:2][0:9])
-
-#pragma omp target exit data map(from: arr[:100])
-
- for (int j = 0; j < 10; j++)
- for (int i = 0; i < 10; i++)
- if (j >= 3 && j < 5 && i < 9)
- assert (arr[j * 10 + i] == 5 * (i + j));
- else
- assert (arr[j * 10 + i] == 0);
-
-
- /* One-dimensional strided update. */
-
- memset (arr, 0, 100 * sizeof (int));
-
-#pragma omp target enter data map(to: arr[:100])
-
- for (int i = 0; i < 100; i++)
- arr[i] = i + 99;
-
-#pragma omp target update to(([100]) arr[3:33:3])
-
-#pragma omp target exit data map(from: arr[:100])
-
- for (int i = 0; i < 100; i++)
- if (i >= 3 && ((i - 3) % 3) == 0)
- assert (arr[i] == i + 99);
- else
- assert (arr[i] == 0);
-
-
- /* One-dimensional strided update without explicit array shape. */
-
- memset (arr, 0, 100 * sizeof (int));
-
-#pragma omp target enter data map(to: arr[:100])
-
- for (int i = 0; i < 100; i++)
- arr[i] = i + 121;
-
-#pragma omp target update to(arr[3:33:3])
-
-#pragma omp target exit data map(from: arr[:100])
-
- for (int i = 0; i < 100; i++)
- if (i >= 3 && ((i - 3) % 3) == 0)
- assert (arr[i] == i + 121);
- else
- assert (arr[i] == 0);
-
- delete[] arr;
-
foo<long> ();
return 0;
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-10.C b/libgomp/testsuite/libgomp.c++/array-shaping-10.C
index 648f02d3479..741042371a0 100644
--- a/libgomp/testsuite/libgomp.c++/array-shaping-10.C
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-10.C
@@ -33,28 +33,6 @@ void foo ()
int main ()
{
- int iarr[N * N];
-
- memset (iarr, 0, N * N * sizeof (int));
-
-#pragma omp target enter data map(to: iarr)
-
-#pragma omp target
- {
- for (int i = 0; i < 10; i++)
- for (int j = 0; j < 10; j++)
- iarr[i * 10 + j] = i + j;
- }
-
- /* An array, but cast to a pointer, then reshaped. */
-#pragma omp target update from(([10][10]) ((int *) &iarr[0])[4:3][4:3])
-
- for (int i = 4; i < 7; i++)
- for (int j = 4; j < 7; j++)
- assert (iarr[i * 10 + j] == i + j);
-
-#pragma omp target exit data map(delete: iarr)
-
foo<unsigned short> ();
return 0;
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-12.C b/libgomp/testsuite/libgomp.c++/array-shaping-12.C
index 103c99aa847..9a924dd3638 100644
--- a/libgomp/testsuite/libgomp.c++/array-shaping-12.C
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-12.C
@@ -35,30 +35,6 @@ void foo ()
int main ()
{
- int iarr_real[N * N];
- int *iarrp = &iarr_real[0];
- int **iarrpp = &iarrp;
-
- memset (iarrp, 0, N * N * sizeof (int));
-
-#pragma omp target enter data map(to: iarr_real)
-
-#pragma omp target
- {
- for (int i = 0; i < 10; i++)
- for (int j = 0; j < 10; j++)
- iarrp[i * 10 + j] = i + j;
- }
-
- /* A pointer with an extra indirection. */
-#pragma omp target update from(([10][10]) (*iarrpp)[4:3][4:3])
-
- for (int i = 4; i < 7; i++)
- for (int j = 4; j < 7; j++)
- assert (iarrp[i * 10 + j] == i + j);
-
-#pragma omp target exit data map(delete: iarr_real)
-
foo<unsigned short> ();
return 0;
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-13.C b/libgomp/testsuite/libgomp.c++/array-shaping-13.C
index 29345ca4264..a5202eb637c 100644
--- a/libgomp/testsuite/libgomp.c++/array-shaping-13.C
+++ b/libgomp/testsuite/libgomp.c++/array-shaping-13.C
@@ -48,41 +48,6 @@ void foo ()
int main ()
{
- int *iptr = new int[N * N * N];
-
- memset (iptr, 0, N * N * N * sizeof (int));
-
-#pragma omp target enter data map(to: iptr[0:N*N*N])
-
-#pragma omp target
- {
- for (int i = 0; i < N; i++)
- for (int j = 0; j < N; j++)
- iptr[i * N * N + 4 * N + j] = i + j;
- }
-
- /* An array ref between two array sections. */
-#pragma omp target update from(([N][N][N]) iptr[2:3][4][6:3])
-
- for (int i = 2; i < 5; i++)
- for (int j = 6; j < 9; j++)
- assert (iptr[i * N * N + 4 * N + j] == i + j);
-
- memset (iptr, 0, N * N * N * sizeof (int));
-
- for (int i = 0; i < N; i++)
- iptr[2 * N * N + i * N + 4] = 3 * i;
-
- /* Array section between two array refs. */
-#pragma omp target update to(([N][N][N]) iptr[2][3:6][4])
-
-#pragma omp target exit data map(from: iptr[0:N*N*N])
-
- for (int i = 3; i < 9; i++)
- assert (iptr[2 * N * N + i * N + 4] == 3 * i);
-
- delete[] iptr;
-
foo<unsigned long> ();
return 0;
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-5.C b/libgomp/testsuite/libgomp.c++/array-shaping-5.C
deleted file mode 100644
index 7046a13c106..00000000000
--- a/libgomp/testsuite/libgomp.c++/array-shaping-5.C
+++ /dev/null
@@ -1,38 +0,0 @@
-// { dg-do run { target offload_device_nonshared_as } }
-
-#include <string.h>
-#include <assert.h>
-
-template<typename T, auto C>
-void foo (T *w, int e, int f, int g)
-{
- memset (w, 0, sizeof (T) * 100);
-
-#pragma omp target enter data map(to: w[:100])
-
- for (int j = 0; j < e; j++)
- for (int i = 0; i < C; i++)
- w[j * C + i] = i + j;
-
-#pragma omp target update to(([e][C]) w[3:2][f:g])
-
-#pragma omp target exit data map(from: w[:100])
-
- for (int j = 0; j < e; j++)
- for (int i = 0; i < C; i++)
- if (j >= 3 && j < 5 && i >= f && i < f + g)
- assert (w[j * C + i] == i + j);
- else
- assert (w[j * C + i] == 0.0f);
-}
-
-int main()
-{
- float *arr = new float[100];
-
- foo<float, 10> (arr, 10, 1, 8);
-
- delete[] arr;
-
- return 0;
-}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/array-shaping-15.c b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-15.c
new file mode 100644
index 00000000000..34867e9e0e0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-15.c
@@ -0,0 +1,41 @@
+/* { dg-do run { target offload_device_nonshared_as } } */
+
+/* The reshaped array is contiguous overall and the section has no explicit
+ stride, but for runtime-variable bounds "f"/"g" the accessed memory may be
+ noncontiguous. */
+
+#include <string.h>
+#include <assert.h>
+
+void
+foo (float *w, int e, int f, int g)
+{
+ memset (w, 0, sizeof (float) * 100);
+
+#pragma omp target enter data map(to: w[:100])
+
+ for (int j = 0; j < e; j++)
+ for (int i = 0; i < 10; i++)
+ w[j * 10 + i] = i + j;
+
+#pragma omp target update to(([e][10]) w[3:2][f:g])
+
+#pragma omp target exit data map(from: w[:100])
+
+ for (int j = 0; j < e; j++)
+ for (int i = 0; i < 10; i++)
+ if (j >= 3 && j < 5 && i >= f && i < f + g)
+ assert (w[j * 10 + i] == i + j);
+ else
+ assert (w[j * 10 + i] == 0.0f);
+}
+
+int
+main (void)
+{
+ float arr[100];
+
+ foo (arr, 10, 1, 8);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-2.C b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-17.c
similarity index 68%
rename from libgomp/testsuite/libgomp.c++/array-shaping-2.C
rename to libgomp/testsuite/libgomp.c-c++-common/array-shaping-17.c
index 027543e8d29..f0dc710b65f 100644
--- a/libgomp/testsuite/libgomp.c++/array-shaping-2.C
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-17.c
@@ -1,12 +1,13 @@
-// { dg-do run { target offload_device_nonshared_as } }
+/* { dg-do run { target offload_device_nonshared_as } } */
+#include <stdlib.h>
#include <string.h>
#include <assert.h>
-template<typename T>
-void foo (T *w)
+void
+foo (int *w)
{
- memset (w, 0, sizeof (T) * 100);
+ memset (w, 0, sizeof (int) * 100);
#pragma omp target enter data map(to: w[:100])
@@ -26,13 +27,14 @@ void foo (T *w)
assert (w[j * 10 + i] == 0);
}
-int main()
+int
+main (void)
{
- int *arr = new int[100];
+ int *arr = (int *) calloc (100, sizeof (int));
- foo<int> (arr);
+ foo (arr);
- delete[] arr;
+ free (arr);
return 0;
}
diff --git a/libgomp/testsuite/libgomp.c/array-shaping-1.c b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-18.c
similarity index 99%
rename from libgomp/testsuite/libgomp.c/array-shaping-1.c
rename to libgomp/testsuite/libgomp.c-c++-common/array-shaping-18.c
index 808c5f9ceae..64e963d4629 100644
--- a/libgomp/testsuite/libgomp.c/array-shaping-1.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-18.c
@@ -12,7 +12,7 @@ int main()
int x = 5;
int arr2d[10][10];
- arr = calloc (100, sizeof (int));
+ arr = (int *) calloc (100, sizeof (int));
/* Update whole reshaped array. */
diff --git a/libgomp/testsuite/libgomp.c/array-shaping-4.c b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-19.c
similarity index 100%
rename from libgomp/testsuite/libgomp.c/array-shaping-4.c
rename to libgomp/testsuite/libgomp.c-c++-common/array-shaping-19.c
diff --git a/libgomp/testsuite/libgomp.c/array-shaping-5.c b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-20.c
similarity index 100%
rename from libgomp/testsuite/libgomp.c/array-shaping-5.c
rename to libgomp/testsuite/libgomp.c-c++-common/array-shaping-20.c
diff --git a/libgomp/testsuite/libgomp.c/array-shaping-6.c b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-21.c
similarity index 94%
rename from libgomp/testsuite/libgomp.c/array-shaping-6.c
rename to libgomp/testsuite/libgomp.c-c++-common/array-shaping-21.c
index 59388232244..753bf0c2605 100644
--- a/libgomp/testsuite/libgomp.c/array-shaping-6.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/array-shaping-21.c
@@ -8,7 +8,7 @@
int main ()
{
- int *iptr = calloc (N * N * N, sizeof (int));
+ int *iptr = (int *) calloc (N * N * N, sizeof (int));
#pragma omp target enter data map(to: iptr[0:N*N*N])
--
2.55.0
From 79141c39aa9c7c9fc5440b39fb23515de2c4bec5 Mon Sep 17 00:00:00 2001
From: Paul-Antoine Arras <[email protected]>
Date: Fri, 18 Sep 2026 14:32:54 +0200
Subject: [PATCH 1/6] WIP Fortran fixups: trailing space + missing hyphen +
constant positive stride check
---
gcc/fortran/openmp.cc | 12 +++++++++++-
.../gfortran.dg/gomp/target-update-iterators-3.f90 | 2 +-
2 files changed, 12 insertions(+), 2 deletions(-)
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 91d42144885..bf687d985bd 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -10200,7 +10200,7 @@ resolve_omp_clauses_aff_dep_map_cache (gfc_code *code,
array isn't contiguous. An expression such as arr(-n:n,-n:n)
could be contiguous even if it looks like it may not be. */
if (code
- && code->op != EXEC_OACC_UPDATE
+ && code->op != EXEC_OACC_UPDATE
&& code->op != EXEC_OMP_TARGET_UPDATE
&& list != OMP_LIST_CACHE
&& list != OMP_LIST_DEPEND
@@ -10260,6 +10260,16 @@ resolve_omp_clauses_aff_dep_map_cache (gfc_code *code,
"in %s clause at %L", name, &n->where);
break;
}
+ else if (ar->stride[i] && code
+ && code->op == EXEC_OMP_TARGET_UPDATE)
+ {
+ if (ar->stride[i]->expr_type == EXPR_CONSTANT
+ && ar->stride[i]->ts.type == BT_INTEGER
+ && 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 (ar->dimen_type[i] != DIMEN_ELEMENT
&& ar->dimen_type[i] != DIMEN_RANGE)
{
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 0ce881b7e40..f7601053ea7 100644
--- a/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90
@@ -23,7 +23,7 @@ end program
! 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 "noncontiguous array section with iterator in target update not supported" 3 "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" } }
--
2.55.0
