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 <tob...@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 <tho...@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 <guoji...@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
signature.asc
Description: PGP signature