[PATCH 3/3] Check array contiguity for OpenACC/Fortran
Tobias Burnus
tobias@codesourcery.com
Tue Jan 28 09:25:00 GMT 2020
(CC: fortran@ as it relates to Fortran.)
Hi all,
On 1/7/20 12:16 PM, Tobias Burnus wrote:
> in terms of the check, it looks fine to me â but I am not sure about
> the spec.
* [OpenACC] Actually, I simply missed the bit (here: OpenACC 3; OpenACC
2.6 is same): âAny array or subarray in a data clause, including Fortran
array pointers, must be a contiguous block of memory, except for dynamic
multidimensional C arrays.â (2.7.1 under Restrictions).
* OpenMP (quoting TR8): âIf a list item is an array section, it must
specify contiguous storage.â (2.22.7.1 map Clause under Restrictions).
However, that one seems to miss the case: ânon_cont_ptr => A(::2); !$omp
... map(non_cont_ptr)â as non_cont_ptr is noncontiguous and an array but
not an array section.
In any case, those are restrictions imposed on the user â which the
compiler may or may not report. (A good one will, I presume. Although,
one can also regard it as implementation defined/vendor extension as GCC
will properly handle those â by mapping also the gaps.)
Cheers,
Tobias
> At least the following test case seems to work fine:
>
> integer :: A(10,10), out(12)
> A = reshape([(i, i=0,100)], shape(A))
> !$omp target map(A(3:6,3:5), out)
> !$acc parallel copy(A(3:6,3:5), out)
> Â out = reshape(A(3:6,3:5), shape(out))
> Â A(3:6,3:5) = 4
> !$acc end parallel
> !$omp end target
> print '(4i3)', out
> print '(10i3)', A
> end
>
> The reason that it works is that the stride is included in the length
> calculation:
> #pragma omp target num_teams(1) thread_limit(0) map(tofrom:MEM[(c_char
> *)_6] [len: _5])
>
> Has for the section (with A(3:6,3:5) -> parm(1:4,1:3)):
> Â Â Â Â Â Â Â Â Â Â Â Â Â parm.1.dim[0].lbound = 1;
> Â Â Â Â Â Â Â Â Â Â Â Â Â parm.1.dim[0].ubound = 4;
> Â Â Â Â Â Â Â Â Â Â Â Â Â parm.1.dim[0].stride = 1;
> Â Â Â Â Â Â Â Â Â Â Â Â Â parm.1.dim[1].lbound = 1;
> Â Â Â Â Â Â Â Â Â Â Â Â Â parm.1.dim[1].ubound = 3;
> Â Â Â Â Â Â Â Â Â Â Â Â Â parm.1.dim[1].stride = 10;
> And instead of doing '(4-1+1) * (3-1+1)' = 12 (i.e. multiplying the
> extends),
> the code does: 'stride * (3-1+1)' = 30.
>
> Cheers,
>
> Tobias
>
> PS: It also works if one puts the stride on the ptr, i.e.:
>
> integer,target :: A(10,10), out(12)
> integer, pointer :: ptr(:,:)
> A = reshape([(i, i=0,100)], shape(A))
> ptr => A(3:6,3:5)
> !$omp target map(ptr, out)
> !$acc parallel copy(ptr, out)
> Â out = reshape(ptr, shape(out))
> Â ptr = 4
> !$acc end parallel
> !$omp end target
> print '(4i3)', out
> print '(10i3)', A
> end
>
> On 1/4/20 3:25 AM, Julian Brown wrote:
>> This patch tightens up error checking for array references used in
>> OpenACC clauses such that they must now be contiguous. I believe this
>> matches up to the spec (as of 2.6). I've tried to make it so an error
>> only triggers if the compiler is sure that a given array expression must
>> be non-contiguous at compile time, although this errs on the side
>> of potentially allowing the user to shoot themselves in the foot at
>> runtime: at least one existing test in the testsuite appears to expect
>> that behaviour.
>> [â¦]
More information about the Gcc-patches
mailing list