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?
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
* * *
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:
(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).
(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.
...
(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.
* * *
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][:][:]
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
* 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.
* The following should be also OK in C/C++ or if simply contiguous in Fortran:
x[i:], x[:i]
as every step is contiguous
* 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.
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)
* If gfc_is_not_contiguous determines that it is noncontiguous
error out (this needs to be added)
* 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)
and a runtime check for
x[2][i][:M][:]
for M == N, unless we always want to reject it?
(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'
* * *
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
