[3/3] Check array contiguity for OpenACC/Fortran

Message ID 20200104022514.115942-3-julian@codesourcery.com
State New
Headers show
Series
  • [1/3] Add OpenACC test for sub-references being pointer or allocatable variables
Related show

Commit Message

Julian Brown Jan. 4, 2020, 2:25 a.m.
Hi,

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.

This hopefully addresses Tobias's concern in:

  https://gcc.gnu.org/ml/gcc-patches/2019-12/msg01439.html

In particular, with the error checking, we no longer get an ICE for the
example given in that message. The new check also catches a test case
that appears to have been relying on undefined behaviour.

Tested with offloading to NVPTX. OK for trunk?

Thanks,

Julian

ChangeLog

2020-01-04  Julian Brown  <julian@codesourcery.com>

	PR fortran/93025

	gcc/fortran/
	* openmp.c (resolve_omp_clauses): Check array references for contiguity.

	gcc/testsuite/
	* gfortran.dg/goacc/mapping-tests-2.f90: New test.
	* gfortran.dg/goacc/subarrays.f95: Expect rejection of non-contiguous
	array.
---
 gcc/fortran/openmp.c                          | 29 +++++++++++++----
 .../gfortran.dg/goacc/mapping-tests-2.f90     | 32 +++++++++++++++++++
 gcc/testsuite/gfortran.dg/goacc/subarrays.f95 |  2 +-
 3 files changed, 55 insertions(+), 8 deletions(-)
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-2.f90

-- 
2.23.0

Comments

Tobias Burnus Jan. 7, 2020, 11:16 a.m. | #1
Hi Julian, hi Thomas,

in terms of the check, it looks fine to me – but I am not sure about the 
spec.

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.

> […]
Tobias Burnus Jan. 28, 2020, 8:50 a.m. | #2
(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.

>> […]

Patch

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 78351b190f7..71308c0235c 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -4533,13 +4533,28 @@  resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 		    /* Look through component refs to find last array
 		       reference.  */
 		    if (openacc && resolved)
-		      while (array_ref
-			     && (array_ref->type == REF_COMPONENT
-				 || (array_ref->type == REF_ARRAY
-				     && array_ref->next
-				     && (array_ref->next->type
-					 == REF_COMPONENT))))
-			array_ref = array_ref->next;
+		      {
+			/* The "!$acc cache" directive allows rectangular
+			   subarrays to be specified, with some restrictions
+			   on the form of bounds (not implemented).
+			   Only raise an error here if we're really sure the
+			   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 (list != OMP_LIST_CACHE
+			    && !gfc_is_simply_contiguous (n->expr, false, true)
+			    && gfc_is_not_contiguous (n->expr))
+			  gfc_error ("Array is not contiguous at %L",
+				     &n->where);
+
+			while (array_ref
+			       && (array_ref->type == REF_COMPONENT
+				   || (array_ref->type == REF_ARRAY
+				       && array_ref->next
+				       && (array_ref->next->type
+					   == REF_COMPONENT))))
+			  array_ref = array_ref->next;
+		      }
 		  }
 		if (array_ref
 		    || (n->expr
diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-2.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-2.f90
new file mode 100644
index 00000000000..1372f6af53e
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-2.f90
@@ -0,0 +1,32 @@ 
+subroutine foo
+  type t
+    integer :: i, j
+  end type t
+
+  type t2
+    type(t) :: cc(3)
+  end type t2
+
+  type(t) x, y(3)
+  type(t2) :: z(3)
+
+  ! OK - map whole aggregated variable
+!$acc enter data copyin(x)
+  ! map(to:x [len: 8])
+
+  ! OK - map two components of the aggregated variable
+!$acc enter data copyin(x%j, x%i)
+
+  ! Bad - we cannot mix full-object and component accesses
+!$acc enter data copyin(x, x%i)
+! { dg-error "Symbol .x. has mixed component and non-component accesses" "" { target "*-*-*" } 21 }
+
+  ! Bad - we cannot do a strided access of 'x'
+  ! No C/C++ equivalent
+!$acc enter data copyin(y(:)%i)
+! { dg-error "Array is not contiguous" "" { target "*-*-*" } 26 }
+
+  ! Bad - again, a strided access
+!$acc enter data copyin(z(1)%cc(:)%i)
+! { dg-error "Array is not contiguous" "" { target "*-*-*" } 30 }
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/subarrays.f95 b/gcc/testsuite/gfortran.dg/goacc/subarrays.f95
index f6adde459f4..fa0378550e9 100644
--- a/gcc/testsuite/gfortran.dg/goacc/subarrays.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/subarrays.f95
@@ -27,7 +27,7 @@  program test
   ! { dg-error "'a' in MAP clause" "" { target *-*-* } .-2 }
   !$acc end parallel
 
-  !$acc parallel copy (b(1:3,2:4))
+  !$acc parallel copy (b(1:3,2:4)) ! { dg-error "Array is not contiguous" }
   !$acc end parallel
   !$acc parallel copy (b(2:3))
   ! { dg-error "Rank mismatch" "" { target *-*-* } .-1 }