Re: [Patch][Fortran] OpenACC – permit common blocks in some clauses

[Thomas Schwinge <thomas at codesourcery dot com>]

Hi Tobias!

By the way, do you know what's the status is for Fortran common blocks in
OpenMP: supported vs. expected per the specification?


On 2019-10-25T16:36:10+0200, Tobias Burnus <tobias@codesourcery.com> wrote:
> On 10/25/19 10:43 AM, Thomas Schwinge wrote:
>> Or, would it be easy to add an OpenACC 'kernels' test case that otherwise
>> faild (at run time, say, with aforementioned duplicate mapping errors, or
>> would contain "strange"/duplicate/conflicting mapping items in the
>> '-fdump-tree-gimple' dump)?

> * I have now a new test case 
> libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90 which looks at 
> omplower.

Thanks.

Curious: why 'omplower' instead of 'gimple' dump?


> Regarding the new test case: Without the gcc/gimplify.c changes, one has 
> (see last item before child fn):
>
>      #pragma omp target oacc_parallel map(tofrom:a [len: 400]) 
> map(tofrom:b [len: 400]) map(tofrom:c [len: 4]) map(tofrom:block [len: 
> 812]) [child fn …
>      #pragma omp target oacc_kernels map(force_tofrom:i [len: 4]) 
> map(tofrom:y [len: 400]) map(tofrom:x [len: 400]) 
> map(tofrom:kernel_block [len: 804]) map(force_tofrom:c [len: 4]) 
> map(tofrom:block [len: 812])  [child fn …
>
> With the changes of gcc/gimplify.c, one has:
>
>      #pragma omp target oacc_parallel map(tofrom:a [len: 400]) 
> map(tofrom:b [len: 400]) map(tofrom:c [len: 4]) [child fn …
>      #pragma omp target oacc_kernels map(force_tofrom:i [len: 4]) 
> map(tofrom:y [len: 400]) map(tofrom:x [len: 400]) map(force_tofrom:c 
> [len: 4])  [child fn …
>
>
> And without gimplify.c, the added run-tests indeed fail with:
> libgomp: Trying to map into device [0x407100..0x407294) object when 
> [0x407100..0x407290) is already mapped

OK, good, my suspicion was thus right that there's something "strange"
there.  ;-)

> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
> @@ -0,0 +1,39 @@
> +! { dg-options "-fopenacc -fdump-tree-omplower" }

(For later: we usually just use 'dg-additional-options
"-fdump-tree-omplower"'; '-fopenacc' is implied inside '*/goacc/'.)

> +
> +module consts
> +  integer, parameter :: n = 100
> +end module consts
> +
> +program main
> +  use consts
> +  implicit none
> +
> +  integer :: i, j
> +  real ::  a(n) = 0, b(n) = 0, c, d
> +  real ::  x(n) = 0, y(n), z
> +  common /BLOCK/ a, b, c, j, d
> +  common /KERNELS_BLOCK/ x, y, z
> +
> +  c = 1.0
> +  !$acc parallel loop copy(/BLOCK/)
> +  do i = 1, n
> +     a(i) = b(i) + c
> +  end do
> +  !$acc kernels
> +  do i = 1, n
> +     x(i) = y(i) + c
> +  end do
> +  !$acc end kernels
> +end program main
> +
> +! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(tofrom:a \\\[len: 400\\\]\\)" 1 "omplower" } }
> +! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(tofrom:b \\\[len: 400\\\]\\\)" 1 "omplower" } }
> +! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } }
> +
> +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\)" 1 "omplower" } }
> +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(tofrom:x \\\[len: 400\\\]\\)" 1 "omplower" } }
> +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(tofrom:y \\\[len: 400\\\]\\\)" 1 "omplower" } }
> +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } }
> +
> +! { dg-final { scan-tree-dump-not "map\\(.*:block\\)" "omplower" } }
> +! { dg-final { scan-tree-dump-not "map\\(.*:kernels_block\\)" "omplower" } }

For my understanding: the several unused variables in the common blocks
are to make sure that they don't cause any issues, don't get mapped at
all?

I we were to add to 'gfortran.dg/goacc/common-block-3.f90' a test case
for the upcoming OpenACC 'serial' construct (which basically equals the
OpenACC 'parallel' construct), would we copy/adapt the 'parallel' 'BLOCK'
test case, or add a new, separate common block?

Or, asking the other way round: why aren't in the current test case,
'parallel' and 'kernels' using the same common block, and both explicitly
'copy' the common block vs. not do that?


> * In the compile-time *{2,3} test case, there is now also a 'enter data' 
> and 'update host/self/device' test.

;-) Heh, 'update' got inside the 'parallel' region:

> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
> @@ -0,0 +1,74 @@
> +[...]
> +  !$acc parallel firstprivate(/blockA/, /blockB/, e, v)
> +  !$acc update device(/blockA/)
> +  !$acc update self(/blockB/, v)
> +  !$acc update host(/blockA/, e, /blockB/)
> +  !$acc end parallel
> +[...]

> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90

Likewise.

As obvoius; see attached, committed "Fix OpenACC directives nesting in
'gfortran.dg/goacc/common-block-1.f90',
'gfortran.dg/goacc/common-block-2.f90'" to trunk in r278047.


Grüße
 Thomas


From 068b41bc6db50d6f600ef95049f41dbbd12f5216 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon, 11 Nov 2019 09:26:40 +0000
Subject: [PATCH] Fix OpenACC directives nesting in
 'gfortran.dg/goacc/common-block-1.f90',
 'gfortran.dg/goacc/common-block-2.f90'

	gcc/testsuite/
	* gfortran.dg/goacc/common-block-1.f90: Fix OpenACC directives
	nesting.
	* gfortran.dg/goacc/common-block-2.f90: Likewise.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@278047 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/testsuite/ChangeLog                            | 6 ++++++
 gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 | 3 ++-
 gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 | 3 ++-
 3 files changed, 10 insertions(+), 2 deletions(-)

diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index cc60856a6a63..f8e626b2fd43 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,9 @@
+2019-11-11  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* gfortran.dg/goacc/common-block-1.f90: Fix OpenACC directives
+	nesting.
+	* gfortran.dg/goacc/common-block-2.f90: Likewise.
+
 2019-11-11  Jiufu Guo  <guojiufu@linux.ibm.com>
 
 	PR tree-optimization/88760
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
index ea437526b464..228637f5883c 100644
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
@@ -55,10 +55,11 @@ program test
   !$acc end parallel
 
   !$acc parallel firstprivate(/blockA/, /blockB/, e, v)
+  !$acc end parallel
+
   !$acc update device(/blockA/)
   !$acc update self(/blockB/, v)
   !$acc update host(/blockA/, e, /blockB/)
-  !$acc end parallel
 
   !$acc enter data pcopyin(/blockA/, /blockB/, e, v)
   !$acc exit data delete(/blockA/, /blockB/, e, v)
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
index 1ba945019f9e..5d49f6195b84 100644
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
@@ -43,10 +43,11 @@ program test
   !$acc end parallel
 
   !$acc parallel firstprivate(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+  !$acc end parallel
+
   !$acc update device(b, /blockA/, x) ! { dg-error "Symbol .x. present on multiple clauses" }
   !$acc update self(z, /blockB/, v) ! { dg-error "Symbol .z. present on multiple clauses" }
   !$acc update host(/blockA/, c) ! { dg-error "Symbol .c. present on multiple clauses" }
-  !$acc end parallel
 
   !$acc enter data copyin(/blockB/, e, v, a, c, y) ! { dg-error "Symbol .y. present on multiple clauses" }
   !$acc exit data delete(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
-- 
2.17.1

Attachment: signature.asc
Description: PGP signature