Re: [PATCH, OpenACC] Fortran "declare create"/allocate support for OpenACC

classic Classic list List threaded Threaded
3 messages Options
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH, OpenACC] Fortran "declare create"/allocate support for OpenACC

Julian Brown-2
On Fri, 21 Sep 2018 03:14:22 +0200
Bernhard Reutner-Fischer <[hidden email]> wrote:

> > diff --git a/gcc/fortran/trans-array.c b/gcc/fortran/trans-array.c
> > index 95ea615..2ac5908 100644
> > --- a/gcc/fortran/trans-array.c
> > +++ b/gcc/fortran/trans-array.c
> > @@ -88,6 +88,7 @@ along with GCC; see the file COPYING3.  If not see
> >  #include "trans-types.h"
> >  #include "trans-array.h"
> >  #include "trans-const.h"
> > +#include "trans-stmt.h"
> >  #include "dependency.h"  
>
> please dont mix declarations and definitions, i.e. please put
> gfc_trans_oacc_declare_allocate() into trans-openmp.c, and add the
> declaration to trans.h, in the corresponding /* In trans-openmp.c */
> block there.
Do you mean like this?

Thanks,

Julian

ChangeLog

2018-09-20  Cesar Philippidis  <[hidden email]>
            Julian Brown  <[hidden email]>

        gcc/
        * omp-low.c (scan_sharing_clauses): Update handling of OpenACC declare
        create, declare copyin and declare deviceptr to have local lifetimes.
        (convert_to_firstprivate_int): Handle pointer types.
        (convert_from_firstprivate_int): Likewise.  Create local storage for
        the values being pointed to.  Add new orig_type argument.
        (lower_omp_target): Handle GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}.
        Add orig_type argument to convert_from_firstprivate_int call.
        Allow pointer types with GOMP_MAP_FIRSTPRIVATE_INT.  Don't privatize
        firstprivate VLAs.
        * tree-pretty-print.c (dump_omp_clause): Handle
        GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}.

        gcc/fortran/
        * gfortran.h (enum gfc_omp_map_op): Add OMP_MAP_DECLARE_ALLOCATE,
        OMP_MAP_DECLARE_DEALLOCATE.
        (gfc_omp_clauses): Add update_allocatable.
        * trans-array.c (gfc_array_allocate): Call
        gfc_trans_oacc_declare_allocate for decls that have oacc_declare_create
        attribute set.
        * trans-decl.c (add_attributes_to_decl): Enable lowering of OpenACC
        declare create, declare copyin and declare deviceptr clauses.
        (add_clause): Don't duplicate OpenACC declare clauses.  Populate
        sym->backend_decl so that it can be used to determine if two symbols are
        unique.
        (find_module_oacc_declare_clauses): Relax oacc_declare_create to
        OMP_MAP_ALLOC, and oacc_declare_copyin to OMP_MAP_TO, in order to
        match OpenACC 2.5 semantics.
        * trans-openmp.c (gfc_trans_omp_clauses): Use GOMP_MAP_ALWAYS_POINTER
        (for update directive) or GOMP_MAP_FIRSTPRIVATE_POINTER (otherwise) for
        allocatable scalar decls.  Handle OMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}
        clauses.
        (gfc_trans_oacc_executable_directive): Use GOMP_MAP_ALWAYS_POINTER
        for allocatable scalar data clauses inside acc update directives.
        (gfc_trans_oacc_declare_allocate): New function.
        * trans-stmt.c (gfc_trans_allocate): Call
        gfc_trans_oacc_declare_allocate for decls with oacc_declare_create
        attribute set.
        (gfc_trans_deallocate): Likewise.
        * trans.h (gfc_trans_oacc_declare_allocate): Declare.

        gcc/testsuite/
        * gfortran.dg/goacc/declare-allocatable-1.f90: New test.

        include/
        * gomp-constants.h (enum gomp_map_kind): Define
        GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE} and GOMP_MAP_FLAG_SPECIAL_4.

        libgomp/
        * oacc-mem.c (gomp_acc_declare_allocate): New function.
        * oacc-parallel.c (GOACC_enter_exit_data): Handle
        GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}.
        * testsuite/libgomp.oacc-fortran/allocatable-array.f90: New test.
        * testsuite/libgomp.oacc-fortran/allocatable-scalar.f90: New test.
        * testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90: New test.
        * testsuite/libgomp.oacc-fortran/declare-allocatable-2.f90: New test.
        * testsuite/libgomp.oacc-fortran/declare-allocatable-3.f90: New test.
        * testsuite/libgomp.oacc-fortran/declare-allocatable-4.f90: New test.

Fortran-declare-create-allocate-support-for-OpenACC-3.patch (36K) Download Attachment
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH, OpenACC] Fortran "declare create"/allocate support for OpenACC

Bernhard Reutner-Fischer
On Sat, 22 Sep 2018 at 00:32, Julian Brown <[hidden email]> wrote:

>
> On Fri, 21 Sep 2018 03:14:22 +0200
> Bernhard Reutner-Fischer <[hidden email]> wrote:
>
> > > diff --git a/gcc/fortran/trans-array.c b/gcc/fortran/trans-array.c
> > > index 95ea615..2ac5908 100644
> > > --- a/gcc/fortran/trans-array.c
> > > +++ b/gcc/fortran/trans-array.c
> > > @@ -88,6 +88,7 @@ along with GCC; see the file COPYING3.  If not see
> > >  #include "trans-types.h"
> > >  #include "trans-array.h"
> > >  #include "trans-const.h"
> > > +#include "trans-stmt.h"
> > >  #include "dependency.h"
> >
> > please dont mix declarations and definitions, i.e. please put
> > gfc_trans_oacc_declare_allocate() into trans-openmp.c, and add the
> > declaration to trans.h, in the corresponding /* In trans-openmp.c */
> > block there.
>
> Do you mean like this?

yes.

@@ -6218,13 +6221,20 @@ add_clause (gfc_symbol *sym, gfc_omp_map_op map_op)
 {
   gfc_omp_namelist *n;

+  if (!module_oacc_clauses)
+    module_oacc_clauses = gfc_get_omp_clauses ();
+
+  if (sym->backend_decl == NULL)
+    gfc_get_symbol_decl (sym);
+
+  for (n = module_oacc_clauses->lists[OMP_LIST_MAP]; n != NULL; n = n->next)
+    if (n->sym->backend_decl == sym->backend_decl)
+      return;
+

Didn't look too close, but should this throw an error instead of
silently returning, or was the error emitted earlier?

Furthermore the testcase uses "call abort" which is non-standard.
We recently moved to "STOP n" in the testsuite, please adjust the new
testcases accordingly.
Like (modulo typos, untested):

$ cat abort_to_stop.awk ; echo EOF
# awk -f ./abort_to_stop.awk < foo.f90 > x && mv x foo.f90
BEGIN { IGNORECASE = 1; i = 1 }
{ while (sub(/call\s\s*abort/, "stop " i)) {let i++;}; print $0; }
EOF

thanks,
Reply | Threaded
Open this post in threaded view
|

[committed] [PR90921] Fortran OpenACC 'declare' directive's module handling causes duplicate data clauses (was: [PATCH, OpenACC] Fortran "declare create"/allocate support for OpenACC)

Thomas Schwinge-8
Hi!

On Thu, 4 Oct 2018 14:04:13 +0100, Julian Brown <[hidden email]> wrote:

> On Sun, 23 Sep 2018 10:48:52 +0200
> Bernhard Reutner-Fischer <[hidden email]> wrote:
>
> > On Sat, 22 Sep 2018 at 00:32, Julian Brown <[hidden email]>
> > wrote:
> >
> > @@ -6218,13 +6221,20 @@ add_clause (gfc_symbol *sym, gfc_omp_map_op
> > map_op) {
> >    gfc_omp_namelist *n;
> >
> > +  if (!module_oacc_clauses)
> > +    module_oacc_clauses = gfc_get_omp_clauses ();
> > +
> > +  if (sym->backend_decl == NULL)
> > +    gfc_get_symbol_decl (sym);
> > +
> > +  for (n = module_oacc_clauses->lists[OMP_LIST_MAP]; n != NULL; n =
> > n->next)
> > +    if (n->sym->backend_decl == sym->backend_decl)
> > +      return;
> > +
> >
> > Didn't look too close, but should this throw an error instead of
> > silently returning, or was the error emitted earlier?
Bernhard, thanks for pointing out this "smelly" code, and then Julian for
analyzing the actual issue:

> The purpose of this fragment seems not to have been to do with error
> reporting at all, but rather to do with de-duplicating symbols that
> are listed (once) in clauses of "declare" directives in module blocks.
> Variables that are listed twice are diagnosed elsewhere.
>
> As for why the de-duplication is necessary, it seems to be because of
> the way that modules are instantiated in programs and in subroutines.
> E.g. in declare-allocatable-1.f90, we have something along the lines of:
>
>   module vars
>     implicit none
>     integer, parameter :: n = 100
>     real*8, allocatable :: b(:)
>    !$acc declare create (b)
>   end module vars
>
>   program test
>     use vars
>     ...
>   end program test
>
>   subroutine sub1
>     use vars
>     ...
>   end subroutine sub1
>
>   subroutine sub2
>     use vars
>     ...
>   end subroutine sub2
>
> The function find_module_oacc_declare_clauses is called for each of
> 'test', 'sub1' and 'sub2'. But in trans-decl.c:finish_oacc_declare, the
> new declare clauses are only attached to the namespace for a FL_PROGRAM
> (i.e. 'test'), not for the subroutines. The module_oacc_clauses global
> variable is reset only after moving the clauses to a FL_PROGRAM's
> namespace, otherwise it accumulates.
>
> Hence, with the above code, we'd scan 'test', find declare clauses, and
> attach them to the namespace for 'test'. We'd then reset
> module_oacc_clauses.
>
> Then, we'd scan 'sub1', and accumulate declare clauses from 'vars' into
> a fresh module_oacc_clauses.
>
> Then we'd scan 'sub2', and accumulate declare clauses from 'vars'
> again: this is why the de-duplication in the patch seemed to be
> necessary.
>
> This seems wrong to me though, and admits the possibility of clauses
> instantiated in a subroutine "leaking" into a subsequent program block.
> As a tentative fix, I've tried resetting module_oacc_clauses before
> each time the find_module_oacc_declare_clauses traversal takes place,
> and removing the de-duplication code.
So, that's clearly a separate bug from everything else discussed as part
of this patch submission; <https://gcc.gnu.org/PR90921> "Fortran OpenACC
'declare' directive's module handling causes duplicate data clauses"
filed.

> This seems to work fine for the current tests in the testsuite, but I
> wonder the reason that things weren't done like like that to start
> with? The code dates back to 2015 (by James Norris):
>
> https://gcc.gnu.org/ml/gcc-patches/2015-11/msg02367.html

There remains a lot of mystery to be resolved regarding the OpenACC
'declare' implementation...  :-(

> --- a/gcc/fortran/trans-decl.c
> +++ b/gcc/fortran/trans-decl.c
> @@ -6272,6 +6275,8 @@ finish_oacc_declare (gfc_namespace *ns, gfc_symbol *sym, bool block)
>    gfc_omp_clauses *omp_clauses = NULL;
>    gfc_omp_namelist *n, *p;
>  
> +  module_oacc_clauses = NULL;
> +
>    gfc_traverse_ns (ns, find_module_oacc_declare_clauses);
>  
>    if (module_oacc_clauses && sym->attr.flavor == FL_PROGRAM)
> @@ -6283,7 +6288,6 @@ finish_oacc_declare (gfc_namespace *ns, gfc_symbol *sym, bool block)
>        new_oc->clauses = module_oacc_clauses;
>  
>        ns->oacc_declare = new_oc;
> -      module_oacc_clauses = NULL;
>      }
>  
>    if (!ns->oacc_declare)
I cannot claim to understand this Fortran OpenACC 'declare' directive's
module handling here, but I can at least confirm via a test case that
I've added, that your change makes the duplicate data clauses go away;
committed to trunk in r272454 "[PR90921] Fortran OpenACC 'declare'
directive's module handling causes duplicate data clauses", see attached.


Grüße
 Thomas



From 9f15ed31065cf6baaae9b3e0e4c16fb9e958fbd9 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Tue, 18 Jun 2019 22:15:53 +0000
Subject: [PATCH] [PR90921] Fortran OpenACC 'declare' directive's module
 handling causes duplicate data clauses

        gcc/fortran/
        PR fortran/90921
        * trans-decl.c (finish_oacc_declare): Reset module_oacc_clauses
        before scanning each namespace.
        gcc/testsuite/
        PR fortran/90921
        * gfortran.dg/goacc/declare-3.f95: Update.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@272454 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/fortran/ChangeLog                         | 6 ++++++
 gcc/fortran/trans-decl.c                      | 2 +-
 gcc/testsuite/ChangeLog                       | 3 +++
 gcc/testsuite/gfortran.dg/goacc/declare-3.f95 | 6 ++++++
 4 files changed, 16 insertions(+), 1 deletion(-)

diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog
index 6fd97b61ce05..32d961ade960 100644
--- a/gcc/fortran/ChangeLog
+++ b/gcc/fortran/ChangeLog
@@ -1,3 +1,9 @@
+2019-06-18  Julian Brown  <[hidden email]>
+
+ PR fortran/90921
+ * trans-decl.c (finish_oacc_declare): Reset module_oacc_clauses
+ before scanning each namespace.
+
 2019-06-18  Thomas Schwinge  <[hidden email]>
 
  PR fortran/85221
diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c
index f504c279c31b..64ce4bba23d9 100644
--- a/gcc/fortran/trans-decl.c
+++ b/gcc/fortran/trans-decl.c
@@ -6491,6 +6491,7 @@ finish_oacc_declare (gfc_namespace *ns, gfc_symbol *sym, bool block)
   gfc_omp_clauses *omp_clauses = NULL;
   gfc_omp_namelist *n, *p;
 
+  module_oacc_clauses = NULL;
   gfc_traverse_ns (ns, find_module_oacc_declare_clauses);
 
   if (module_oacc_clauses && sym->attr.flavor == FL_PROGRAM)
@@ -6502,7 +6503,6 @@ finish_oacc_declare (gfc_namespace *ns, gfc_symbol *sym, bool block)
       new_oc->clauses = module_oacc_clauses;
 
       ns->oacc_declare = new_oc;
-      module_oacc_clauses = NULL;
     }
 
   if (!ns->oacc_declare)
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 552ccc6fbd68..6ff197c8e4dc 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,5 +1,8 @@
 2019-06-18  Thomas Schwinge  <[hidden email]>
 
+ PR fortran/90921
+ * gfortran.dg/goacc/declare-3.f95: Update.
+
  PR fortran/85221
  * gfortran.dg/goacc/declare-3.f95: New file.
 
diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-3.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-3.f95
index ec5d4c5a062a..80d9903a9dc6 100644
--- a/gcc/testsuite/gfortran.dg/goacc/declare-3.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/declare-3.f95
@@ -1,5 +1,7 @@
 ! Test valid usage of the OpenACC 'declare' directive.
 
+! { dg-additional-options "-fdump-tree-original" }
+
 module mod_a
   implicit none
   integer :: a
@@ -44,4 +46,8 @@ program test
   use mod_c
   use mod_d
   use mod_e
+
+  ! { dg-final { scan-tree-dump {(?n)#pragma acc data map\(force_alloc:d\) map\(force_deviceptr:c\) map\(force_to:b\) map\(force_alloc:a\)$} original } }
 end program test
+
+! { dg-final { scan-tree-dump-times {#pragma acc data} 1 original } }
--
2.20.1


signature.asc (671 bytes) Download Attachment