[PATCH] Fix component mappings with derived types for OpenACC

[PATCH] Fix component mappings with derived types for OpenACC

[Julian Brown]

[-- Attachment #1: Type: text/plain, Size: 1317 bytes --]

Hi,

This patch fixes a bug with mapping Fortran components (i.e. with the
manual deep-copy support) which themselves have derived types. I've
also added a couple of new tests to make sure such mappings are lowered
correctly, and to check for the case that Tobias found in the message:

  https://gcc.gnu.org/ml/gcc-patches/2020-01/msg00215.html

The previous incorrect mapping was causing the error condition mentioned
in that message to fail to trigger, and I think (my!) code in libgomp
(goacc_exit_data_internal) to handle GOMP_MAP_STRUCT specially was
papering over the bad mapping also. Oops!

I haven't attempted to implement the (harder) sub-copy detection, if
that is indeed supposed to be forbidden by the spec. This patch should
get us to the same behaviour in Fortran as in C & C++ though.

Tested with offloading to nvptx, also with the (uncommitted)
reference-count self-checking patch enabled.

OK?

Thanks,

Julian

ChangeLog

    gcc/fortran/
    * trans-openmp.c (gfc_trans_omp_clauses): Use inner not decl for
    components with derived types.

    gcc/testsuite/
    * gfortran.dg/goacc/mapping-tests-3.f90: New test.
    * gfortran.dg/goacc/mapping-tests-4.f90: New test.

    libgomp/
    * oacc-mem.c (goacc_exit_data_internal): Remove special (no-copyback)
    behaviour for GOMP_MAP_STRUCT.

[-- Attachment #2: component-mappings-derived-types-1.diff --]
[-- Type: text/x-patch, Size: 3656 bytes --]

commit 5e9d8846bbaa33a9bdb08adcf1ee9f224a8e8fc0
Author: Julian Brown <julian@codesourcery.com>
Date:   Wed Jan 8 15:57:46 2020 -0800

    Fix component mappings with derived types for OpenACC
    
            gcc/fortran/
            * trans-openmp.c (gfc_trans_omp_clauses): Use inner not decl for
            components with derived types.
    
            gcc/testsuite/
            * gfortran.dg/goacc/mapping-tests-3.f90: New test.
            * gfortran.dg/goacc/mapping-tests-4.f90: New test.
    
            libgomp/
            * oacc-mem.c (goacc_exit_data_internal): Remove special (no-copyback)
            behaviour for GOMP_MAP_STRUCT.

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 9835d2aae3c..efc392d7fa6 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2783,9 +2783,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			}
 		      else
 			{
-			  OMP_CLAUSE_DECL (node) = decl;
+			  OMP_CLAUSE_DECL (node) = inner;
 			  OMP_CLAUSE_SIZE (node)
-			    = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+			    = TYPE_SIZE_UNIT (TREE_TYPE (inner));
 			}
 		    }
 		  else if (lastcomp->next
diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
new file mode 100644
index 00000000000..312f596461e
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
@@ -0,0 +1,15 @@
+! { dg-options "-fopenacc -fdump-tree-omplower" }
+
+subroutine foo
+  type one
+    integer i, j
+  end type
+  type two
+    type(one) A, B
+  end type
+
+  type(two) x
+
+  !$acc enter data copyin(x%A)
+! { dg-final { scan-tree-dump-times "omp target oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\)" 1 "omplower" } }
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
new file mode 100644
index 00000000000..6257af942df
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
@@ -0,0 +1,17 @@
+subroutine foo
+  type one
+    integer i, j
+  end type
+  type two
+    type(one) A, B
+  end type
+
+  type(two) x
+
+! This is accepted at present, although it represents a probably-unintentional
+! overlapping subcopy.
+  !$acc enter data copyin(x%A, x%A%i)
+! But this raises an error.
+  !$acc enter data copyin(x%A, x%A%i, x%A%i)
+! { dg-error ".x.a.i. appears more than once in map clauses" "" { target "*-*-*" } 15 }
+end
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 2d4bba78efd..232683a85f0 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1136,38 +1136,6 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 	  break;
 
 	case GOMP_MAP_STRUCT:
-	  {
-	    int elems = sizes[i];
-	    for (int j = 1; j <= elems; j++)
-	      {
-		struct splay_tree_key_s k;
-		k.host_start = (uintptr_t) hostaddrs[i + j];
-		k.host_end = k.host_start + sizes[i + j];
-		splay_tree_key str;
-		str = splay_tree_lookup (&acc_dev->mem_map, &k);
-		if (str)
-		  {
-		    if (finalize)
-		      {
-			if (str->refcount != REFCOUNT_INFINITY)
-			  str->refcount -= str->virtual_refcount;
-			str->virtual_refcount = 0;
-		      }
-		    if (str->virtual_refcount > 0)
-		      {
-			if (str->refcount != REFCOUNT_INFINITY)
-			  str->refcount--;
-			str->virtual_refcount--;
-		      }
-		    else if (str->refcount > 0
-			     && str->refcount != REFCOUNT_INFINITY)
-		      str->refcount--;
-		    if (str->refcount == 0)
-		      gomp_remove_var_async (acc_dev, str, aq);
-		  }
-	      }
-	    i += elems;
-	  }
 	  break;
 
 	default:

Re: [PATCH] Fix component mappings with derived types for OpenACC

[Tobias Burnus]

Hi Julian,

the gfortran part is rather obvious and it and the test case look fine 
to me → OK.
The oacc-mem.c also looks okay, but I assume Thomas needs to 
rubber-stamp it.

Tobias

On 1/10/20 2:49 AM, Julian Brown wrote:
> This patch fixes a bug with mapping Fortran components (i.e. with the
> manual deep-copy support) which themselves have derived types. I've
> also added a couple of new tests to make sure such mappings are lowered
> correctly, and to check for the case that Tobias found in the message:
>
>    https://gcc.gnu.org/ml/gcc-patches/2020-01/msg00215.html
>
> The previous incorrect mapping was causing the error condition mentioned
> in that message to fail to trigger, and I think (my!) code in libgomp
> (goacc_exit_data_internal) to handle GOMP_MAP_STRUCT specially was
> papering over the bad mapping also. Oops!
>
> I haven't attempted to implement the (harder) sub-copy detection, if
> that is indeed supposed to be forbidden by the spec. This patch should
> get us to the same behaviour in Fortran as in C & C++ though.
>
> Tested with offloading to nvptx, also with the (uncommitted)
> reference-count self-checking patch enabled.
>
> OK?
>
> Thanks,
>
> Julian
>
> ChangeLog
>
>      gcc/fortran/
>      * trans-openmp.c (gfc_trans_omp_clauses): Use inner not decl for
>      components with derived types.
>
>      gcc/testsuite/
>      * gfortran.dg/goacc/mapping-tests-3.f90: New test.
>      * gfortran.dg/goacc/mapping-tests-4.f90: New test.
>
>      libgomp/
>      * oacc-mem.c (goacc_exit_data_internal): Remove special (no-copyback)
>      behaviour for GOMP_MAP_STRUCT.
>
> component-mappings-derived-types-1.diff
>
> commit 5e9d8846bbaa33a9bdb08adcf1ee9f224a8e8fc0
> Author: Julian Brown<julian@codesourcery.com>
> Date:   Wed Jan 8 15:57:46 2020 -0800
>
>      Fix component mappings with derived types for OpenACC
>      
>              gcc/fortran/
>              * trans-openmp.c (gfc_trans_omp_clauses): Use inner not decl for
>              components with derived types.
>      
>              gcc/testsuite/
>              * gfortran.dg/goacc/mapping-tests-3.f90: New test.
>              * gfortran.dg/goacc/mapping-tests-4.f90: New test.
>      
>              libgomp/
>              * oacc-mem.c (goacc_exit_data_internal): Remove special (no-copyback)
>              behaviour for GOMP_MAP_STRUCT.
>
> diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
> index 9835d2aae3c..efc392d7fa6 100644
> --- a/gcc/fortran/trans-openmp.c
> +++ b/gcc/fortran/trans-openmp.c
> @@ -2783,9 +2783,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>   			}
>   		      else
>   			{
> -			  OMP_CLAUSE_DECL (node) = decl;
> +			  OMP_CLAUSE_DECL (node) = inner;
>   			  OMP_CLAUSE_SIZE (node)
> -			    = TYPE_SIZE_UNIT (TREE_TYPE (decl));
> +			    = TYPE_SIZE_UNIT (TREE_TYPE (inner));
>   			}
>   		    }
>   		  else if (lastcomp->next
> diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
> new file mode 100644
> index 00000000000..312f596461e
> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
> @@ -0,0 +1,15 @@
> +! { dg-options "-fopenacc -fdump-tree-omplower" }
> +
> +subroutine foo
> +  type one
> +    integer i, j
> +  end type
> +  type two
> +    type(one) A, B
> +  end type
> +
> +  type(two) x
> +
> +  !$acc enter data copyin(x%A)
> +! { dg-final { scan-tree-dump-times "omp target oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\)" 1 "omplower" } }
> +end
> diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
> new file mode 100644
> index 00000000000..6257af942df
> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
> @@ -0,0 +1,17 @@
> +subroutine foo
> +  type one
> +    integer i, j
> +  end type
> +  type two
> +    type(one) A, B
> +  end type
> +
> +  type(two) x
> +
> +! This is accepted at present, although it represents a probably-unintentional
> +! overlapping subcopy.
> +  !$acc enter data copyin(x%A, x%A%i)
> +! But this raises an error.
> +  !$acc enter data copyin(x%A, x%A%i, x%A%i)
> +! { dg-error ".x.a.i. appears more than once in map clauses" "" { target "*-*-*" } 15 }
> +end
> diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
> index 2d4bba78efd..232683a85f0 100644
> --- a/libgomp/oacc-mem.c
> +++ b/libgomp/oacc-mem.c
> @@ -1136,38 +1136,6 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
>   	  break;
>   
>   	case GOMP_MAP_STRUCT:
> -	  {
> -	    int elems = sizes[i];
> -	    for (int j = 1; j <= elems; j++)
> -	      {
> -		struct splay_tree_key_s k;
> -		k.host_start = (uintptr_t) hostaddrs[i + j];
> -		k.host_end = k.host_start + sizes[i + j];
> -		splay_tree_key str;
> -		str = splay_tree_lookup (&acc_dev->mem_map, &k);
> -		if (str)
> -		  {
> -		    if (finalize)
> -		      {
> -			if (str->refcount != REFCOUNT_INFINITY)
> -			  str->refcount -= str->virtual_refcount;
> -			str->virtual_refcount = 0;
> -		      }
> -		    if (str->virtual_refcount > 0)
> -		      {
> -			if (str->refcount != REFCOUNT_INFINITY)
> -			  str->refcount--;
> -			str->virtual_refcount--;
> -		      }
> -		    else if (str->refcount > 0
> -			     && str->refcount != REFCOUNT_INFINITY)
> -		      str->refcount--;
> -		    if (str->refcount == 0)
> -		      gomp_remove_var_async (acc_dev, str, aq);
> -		  }
> -	      }
> -	    i += elems;
> -	  }
>   	  break;
>   
>   	default:

Re: [PATCH] Fix component mappings with derived types for OpenACC

[Julian Brown]

On Fri, 24 Jan 2020 10:58:49 +0100
Tobias Burnus <tobias@codesourcery.com> wrote:

> Hi Julian,
> 
> the gfortran part is rather obvious and it and the test case look
> fine to me → OK.
> The oacc-mem.c also looks okay, but I assume Thomas needs to 
> rubber-stamp it.

I understand that Thomas is unavailable for the time being, so won't be
able to use his rubber-stamp powers. I added the offending libgomp code
to start with though, so I think I can go ahead and commit the patch.
I'll hold off for 24 hours though in case there are any objections
(Jakub?).

Thanks,

Julian

> On 1/10/20 2:49 AM, Julian Brown wrote:
> > This patch fixes a bug with mapping Fortran components (i.e. with
> > the manual deep-copy support) which themselves have derived types.
> > I've also added a couple of new tests to make sure such mappings
> > are lowered correctly, and to check for the case that Tobias found
> > in the message:
> >
> >    https://gcc.gnu.org/ml/gcc-patches/2020-01/msg00215.html
> >
> > The previous incorrect mapping was causing the error condition
> > mentioned in that message to fail to trigger, and I think (my!)
> > code in libgomp (goacc_exit_data_internal) to handle
> > GOMP_MAP_STRUCT specially was papering over the bad mapping also.
> > Oops!
> >
> > I haven't attempted to implement the (harder) sub-copy detection, if
> > that is indeed supposed to be forbidden by the spec. This patch
> > should get us to the same behaviour in Fortran as in C & C++ though.
> >
> > Tested with offloading to nvptx, also with the (uncommitted)
> > reference-count self-checking patch enabled.
> >
> > OK?
> >
> > Thanks,
> >
> > Julian
> >
> > ChangeLog
> >
> >      gcc/fortran/
> >      * trans-openmp.c (gfc_trans_omp_clauses): Use inner not decl
> > for components with derived types.
> >
> >      gcc/testsuite/
> >      * gfortran.dg/goacc/mapping-tests-3.f90: New test.
> >      * gfortran.dg/goacc/mapping-tests-4.f90: New test.
> >
> >      libgomp/
> >      * oacc-mem.c (goacc_exit_data_internal): Remove special
> > (no-copyback) behaviour for GOMP_MAP_STRUCT.
> >
> > component-mappings-derived-types-1.diff
> >
> > commit 5e9d8846bbaa33a9bdb08adcf1ee9f224a8e8fc0
> > Author: Julian Brown<julian@codesourcery.com>
> > Date:   Wed Jan 8 15:57:46 2020 -0800
> >
> >      Fix component mappings with derived types for OpenACC
> >      
> >              gcc/fortran/
> >              * trans-openmp.c (gfc_trans_omp_clauses): Use inner
> > not decl for components with derived types.
> >      
> >              gcc/testsuite/
> >              * gfortran.dg/goacc/mapping-tests-3.f90: New test.
> >              * gfortran.dg/goacc/mapping-tests-4.f90: New test.
> >      
> >              libgomp/
> >              * oacc-mem.c (goacc_exit_data_internal): Remove
> > special (no-copyback) behaviour for GOMP_MAP_STRUCT.
> >
> > diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
> > index 9835d2aae3c..efc392d7fa6 100644
> > --- a/gcc/fortran/trans-openmp.c
> > +++ b/gcc/fortran/trans-openmp.c
> > @@ -2783,9 +2783,9 @@ gfc_trans_omp_clauses (stmtblock_t *block,
> > gfc_omp_clauses *clauses, }
> >   		      else
> >   			{
> > -			  OMP_CLAUSE_DECL (node) = decl;
> > +			  OMP_CLAUSE_DECL (node) = inner;
> >   			  OMP_CLAUSE_SIZE (node)
> > -			    = TYPE_SIZE_UNIT (TREE_TYPE (decl));
> > +			    = TYPE_SIZE_UNIT (TREE_TYPE (inner));
> >   			}
> >   		    }
> >   		  else if (lastcomp->next
> > diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
> > b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 new file mode
> > 100644 index 00000000000..312f596461e
> > --- /dev/null
> > +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
> > @@ -0,0 +1,15 @@
> > +! { dg-options "-fopenacc -fdump-tree-omplower" }
> > +
> > +subroutine foo
> > +  type one
> > +    integer i, j
> > +  end type
> > +  type two
> > +    type(one) A, B
> > +  end type
> > +
> > +  type(two) x
> > +
> > +  !$acc enter data copyin(x%A)
> > +! { dg-final { scan-tree-dump-times "omp target
> > oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a
> > \\\[len: \[0-9\]+\\\]\\)" 1 "omplower" } } +end diff --git
> > a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
> > b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 new file mode
> > 100644 index 00000000000..6257af942df --- /dev/null
> > +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
> > @@ -0,0 +1,17 @@
> > +subroutine foo
> > +  type one
> > +    integer i, j
> > +  end type
> > +  type two
> > +    type(one) A, B
> > +  end type
> > +
> > +  type(two) x
> > +
> > +! This is accepted at present, although it represents a
> > probably-unintentional +! overlapping subcopy.
> > +  !$acc enter data copyin(x%A, x%A%i)
> > +! But this raises an error.
> > +  !$acc enter data copyin(x%A, x%A%i, x%A%i)
> > +! { dg-error ".x.a.i. appears more than once in map clauses"
> > "" { target "*-*-*" } 15 } +end
> > diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
> > index 2d4bba78efd..232683a85f0 100644
> > --- a/libgomp/oacc-mem.c
> > +++ b/libgomp/oacc-mem.c
> > @@ -1136,38 +1136,6 @@ goacc_exit_data_internal (struct
> > gomp_device_descr *acc_dev, size_t mapnum, break;
> >   
> >   	case GOMP_MAP_STRUCT:
> > -	  {
> > -	    int elems = sizes[i];
> > -	    for (int j = 1; j <= elems; j++)
> > -	      {
> > -		struct splay_tree_key_s k;
> > -		k.host_start = (uintptr_t) hostaddrs[i + j];
> > -		k.host_end = k.host_start + sizes[i + j];
> > -		splay_tree_key str;
> > -		str = splay_tree_lookup (&acc_dev->mem_map, &k);
> > -		if (str)
> > -		  {
> > -		    if (finalize)
> > -		      {
> > -			if (str->refcount != REFCOUNT_INFINITY)
> > -			  str->refcount -= str->virtual_refcount;
> > -			str->virtual_refcount = 0;
> > -		      }
> > -		    if (str->virtual_refcount > 0)
> > -		      {
> > -			if (str->refcount != REFCOUNT_INFINITY)
> > -			  str->refcount--;
> > -			str->virtual_refcount--;
> > -		      }
> > -		    else if (str->refcount > 0
> > -			     && str->refcount != REFCOUNT_INFINITY)
> > -		      str->refcount--;
> > -		    if (str->refcount == 0)
> > -		      gomp_remove_var_async (acc_dev, str, aq);
> > -		  }
> > -	      }
> > -	    i += elems;
> > -	  }
> >   	  break;
> >   
> >   	default:  

Re: [PATCH] Fix component mappings with derived types for OpenACC

[Thomas Schwinge]

Hi!

On 2020-01-28T13:41:00+0000, Julian Brown <julian@codesourcery.com> wrote:
> On Fri, 24 Jan 2020 10:58:49 +0100
> Tobias Burnus <tobias@codesourcery.com> wrote:
>> the gfortran part is rather obvious and it and the test case look
>> fine to me → OK.
>> The oacc-mem.c also looks okay, but I assume Thomas needs to
>> rubber-stamp it.
>
> I understand that Thomas is unavailable for the time being, so won't be
> able to use his rubber-stamp powers. I added the offending libgomp code
> to start with though, so I think I can go ahead and commit the patch.
> I'll hold off for 24 hours though in case there are any objections
> (Jakub?).

So, in the end you didn't commit this, and now we've got the "un-fixed"
OpenACC/Fortran code generation in the GCC 10.1 release, so have to deal
with it in one way or another going forward, regarding libgomp ABI
compatibility.  (Ideally), we need to make sure that "un-fixed" GCC
10.1-built executables continue to work "as good as before" when
dynamically linked/running against "fixed" GCC 10.1+ shared libraries.

Do we get such desired behavior by the patch quoted below?  In
particular: removing the 'GOMP_MAP_STRUCT' handling code, but leaving in
the empty 'case GOMP_MAP_STRUCT:' as a no-op, so that we don't run into
'default:' case 'goacc_exit_data_internal UNHANDLED kind'?  Is that
sufficient?

Is my understanding correct that "fixed" GCC won't generate such
'GOMP_MAP_STRUCT' anymore (I have't studied in detail), and this empty
'case GOMP_MAP_STRUCT:' only remains in here for backwards compatibility?
In this case, please add a comment to the code, stating this.  Otherwise,
please add a comment why "do nothing" is appropriate for
'GOMP_MAP_STRUCT'.  In particular, for both scenarios, why we don't need
to skip the following 'sizes[i]' mappings?


Grüße
 Thomas


>> On 1/10/20 2:49 AM, Julian Brown wrote:
>> > This patch fixes a bug with mapping Fortran components (i.e. with
>> > the manual deep-copy support) which themselves have derived types.
>> > I've also added a couple of new tests to make sure such mappings
>> > are lowered correctly, and to check for the case that Tobias found
>> > in the message:
>> >
>> >    https://gcc.gnu.org/ml/gcc-patches/2020-01/msg00215.html
>> >
>> > The previous incorrect mapping was causing the error condition
>> > mentioned in that message to fail to trigger, and I think (my!)
>> > code in libgomp (goacc_exit_data_internal) to handle
>> > GOMP_MAP_STRUCT specially was papering over the bad mapping also.
>> > Oops!
>> >
>> > I haven't attempted to implement the (harder) sub-copy detection, if
>> > that is indeed supposed to be forbidden by the spec. This patch
>> > should get us to the same behaviour in Fortran as in C & C++ though.
>> >
>> > Tested with offloading to nvptx, also with the (uncommitted)
>> > reference-count self-checking patch enabled.
>> >
>> > OK?
>> >
>> > Thanks,
>> >
>> > Julian
>> >
>> > ChangeLog
>> >
>> >      gcc/fortran/
>> >      * trans-openmp.c (gfc_trans_omp_clauses): Use inner not decl
>> > for components with derived types.
>> >
>> >      gcc/testsuite/
>> >      * gfortran.dg/goacc/mapping-tests-3.f90: New test.
>> >      * gfortran.dg/goacc/mapping-tests-4.f90: New test.
>> >
>> >      libgomp/
>> >      * oacc-mem.c (goacc_exit_data_internal): Remove special
>> > (no-copyback) behaviour for GOMP_MAP_STRUCT.
>> >
>> > component-mappings-derived-types-1.diff
>> >
>> > commit 5e9d8846bbaa33a9bdb08adcf1ee9f224a8e8fc0
>> > Author: Julian Brown<julian@codesourcery.com>
>> > Date:   Wed Jan 8 15:57:46 2020 -0800
>> >
>> >      Fix component mappings with derived types for OpenACC
>> >
>> >              gcc/fortran/
>> >              * trans-openmp.c (gfc_trans_omp_clauses): Use inner
>> > not decl for components with derived types.
>> >
>> >              gcc/testsuite/
>> >              * gfortran.dg/goacc/mapping-tests-3.f90: New test.
>> >              * gfortran.dg/goacc/mapping-tests-4.f90: New test.
>> >
>> >              libgomp/
>> >              * oacc-mem.c (goacc_exit_data_internal): Remove
>> > special (no-copyback) behaviour for GOMP_MAP_STRUCT.
>> >
>> > diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
>> > index 9835d2aae3c..efc392d7fa6 100644
>> > --- a/gcc/fortran/trans-openmp.c
>> > +++ b/gcc/fortran/trans-openmp.c
>> > @@ -2783,9 +2783,9 @@ gfc_trans_omp_clauses (stmtblock_t *block,
>> > gfc_omp_clauses *clauses, }
>> >                  else
>> >                    {
>> > -                    OMP_CLAUSE_DECL (node) = decl;
>> > +                    OMP_CLAUSE_DECL (node) = inner;
>> >                      OMP_CLAUSE_SIZE (node)
>> > -                      = TYPE_SIZE_UNIT (TREE_TYPE (decl));
>> > +                      = TYPE_SIZE_UNIT (TREE_TYPE (inner));
>> >                    }
>> >                }
>> >              else if (lastcomp->next
>> > diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
>> > b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 new file mode
>> > 100644 index 00000000000..312f596461e
>> > --- /dev/null
>> > +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
>> > @@ -0,0 +1,15 @@
>> > +! { dg-options "-fopenacc -fdump-tree-omplower" }
>> > +
>> > +subroutine foo
>> > +  type one
>> > +    integer i, j
>> > +  end type
>> > +  type two
>> > +    type(one) A, B
>> > +  end type
>> > +
>> > +  type(two) x
>> > +
>> > +  !$acc enter data copyin(x%A)
>> > +! { dg-final { scan-tree-dump-times "omp target
>> > oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a
>> > \\\[len: \[0-9\]+\\\]\\)" 1 "omplower" } } +end diff --git
>> > a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
>> > b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 new file mode
>> > 100644 index 00000000000..6257af942df --- /dev/null
>> > +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
>> > @@ -0,0 +1,17 @@
>> > +subroutine foo
>> > +  type one
>> > +    integer i, j
>> > +  end type
>> > +  type two
>> > +    type(one) A, B
>> > +  end type
>> > +
>> > +  type(two) x
>> > +
>> > +! This is accepted at present, although it represents a
>> > probably-unintentional +! overlapping subcopy.
>> > +  !$acc enter data copyin(x%A, x%A%i)
>> > +! But this raises an error.
>> > +  !$acc enter data copyin(x%A, x%A%i, x%A%i)
>> > +! { dg-error ".x.a.i. appears more than once in map clauses"
>> > "" { target "*-*-*" } 15 } +end
>> > diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
>> > index 2d4bba78efd..232683a85f0 100644
>> > --- a/libgomp/oacc-mem.c
>> > +++ b/libgomp/oacc-mem.c
>> > @@ -1136,38 +1136,6 @@ goacc_exit_data_internal (struct
>> > gomp_device_descr *acc_dev, size_t mapnum, break;
>> >
>> >    case GOMP_MAP_STRUCT:
>> > -    {
>> > -      int elems = sizes[i];
>> > -      for (int j = 1; j <= elems; j++)
>> > -        {
>> > -          struct splay_tree_key_s k;
>> > -          k.host_start = (uintptr_t) hostaddrs[i + j];
>> > -          k.host_end = k.host_start + sizes[i + j];
>> > -          splay_tree_key str;
>> > -          str = splay_tree_lookup (&acc_dev->mem_map, &k);
>> > -          if (str)
>> > -            {
>> > -              if (finalize)
>> > -                {
>> > -                  if (str->refcount != REFCOUNT_INFINITY)
>> > -                    str->refcount -= str->virtual_refcount;
>> > -                  str->virtual_refcount = 0;
>> > -                }
>> > -              if (str->virtual_refcount > 0)
>> > -                {
>> > -                  if (str->refcount != REFCOUNT_INFINITY)
>> > -                    str->refcount--;
>> > -                  str->virtual_refcount--;
>> > -                }
>> > -              else if (str->refcount > 0
>> > -                       && str->refcount != REFCOUNT_INFINITY)
>> > -                str->refcount--;
>> > -              if (str->refcount == 0)
>> > -                gomp_remove_var_async (acc_dev, str, aq);
>> > -            }
>> > -        }
>> > -      i += elems;
>> > -    }
>> >      break;
>> >
>> >    default:
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

[PATCH 0/3] OpenACC "exit data" copyout, and Fortran derived-type members

[Julian Brown]

This is a split version of the patch previously posted here:

  https://gcc.gnu.org/legacy-ml/gcc-patches/2020-01/msg00512.html

There were actually a couple of different issues addressed by that patch,
and Thomas raised another (or so).  This separated-out version should
be more-obviously safe (in terms of maintaining backwards compatibility,
where that makes sense).  Further commentary on individual patches.

Tested (as a series) with offloading to NVPTX. OK?

Julian

Julian Brown (3):
  OpenACC "exit data" copyout for struct members
  Strip GOMP_MAP_STRUCT from OpenACC exit data mappings
  Fortran derived-type mapping fix

 gcc/fortran/trans-openmp.c                    |  4 +-
 gcc/gimplify.c                                |  3 +-
 .../goacc/struct-enter-exit-data-1.c          | 35 +++++++++++++++
 .../gfortran.dg/goacc/mapping-tests-3.f90     | 15 +++++++
 .../gfortran.dg/goacc/mapping-tests-4.f90     | 17 +++++++
 libgomp/oacc-mem.c                            | 32 --------------
 .../struct-copyout-1.c                        | 38 ++++++++++++++++
 .../struct-copyout-2.c                        | 44 +++++++++++++++++++
 8 files changed, 153 insertions(+), 35 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c

-- 
2.23.0

[PATCH 1/3] OpenACC "exit data" copyout for struct members

[Julian Brown]

This patch removes unnecessary special-case code in oacc-mem.c
which prevented copy-back of structure members on OpenACC "exit data"
directives.  I've added a couple of new testcases to verify the behaviour
(which fail without the patch).

(On editing this mail, I notice I omitted to add a comment about
GOMP_MAP_STRUCT being a no-op, and/or no longer emitted for OpenACC
exit data with the patch later in this series.  I'll add such a comment
before commit.)

OK?

Thanks,

Julian

	libgomp/
	* oacc-mem.c (goacc_exit_data_internal): Remove special-case
	(no-copyback) handling for GOMP_MAP_STRUCT.
	* testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c: New test.
---
 libgomp/oacc-mem.c                            | 32 --------------
 .../struct-copyout-1.c                        | 38 ++++++++++++++++
 .../struct-copyout-2.c                        | 44 +++++++++++++++++++
 3 files changed, 82 insertions(+), 32 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 2d4bba78efd..232683a85f0 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1136,38 +1136,6 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 	  break;
 
 	case GOMP_MAP_STRUCT:
-	  {
-	    int elems = sizes[i];
-	    for (int j = 1; j <= elems; j++)
-	      {
-		struct splay_tree_key_s k;
-		k.host_start = (uintptr_t) hostaddrs[i + j];
-		k.host_end = k.host_start + sizes[i + j];
-		splay_tree_key str;
-		str = splay_tree_lookup (&acc_dev->mem_map, &k);
-		if (str)
-		  {
-		    if (finalize)
-		      {
-			if (str->refcount != REFCOUNT_INFINITY)
-			  str->refcount -= str->virtual_refcount;
-			str->virtual_refcount = 0;
-		      }
-		    if (str->virtual_refcount > 0)
-		      {
-			if (str->refcount != REFCOUNT_INFINITY)
-			  str->refcount--;
-			str->virtual_refcount--;
-		      }
-		    else if (str->refcount > 0
-			     && str->refcount != REFCOUNT_INFINITY)
-		      str->refcount--;
-		    if (str->refcount == 0)
-		      gomp_remove_var_async (acc_dev, str, aq);
-		  }
-	      }
-	    i += elems;
-	  }
 	  break;
 
 	default:
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
new file mode 100644
index 00000000000..b86f1c921a9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
@@ -0,0 +1,38 @@
+#include <assert.h>
+
+struct str1 {
+  int a;
+  int b;
+};
+
+struct str2 {
+  int c;
+  int d;
+  struct str1 s;
+};
+
+int
+main (int argc, char *argv[])
+{
+  struct str2 t;
+
+  t.c = 1;
+  t.d = 2;
+  t.s.a = 3;
+  t.s.b = 4;
+
+  #pragma acc enter data copyin(t.s)
+
+  #pragma acc serial present(t.s) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
+  {
+    t.s.a = 5;
+    t.s.b = 6;
+  }
+
+  #pragma acc exit data copyout(t.s)
+
+  assert (t.s.a == 5);
+  assert (t.s.b == 6);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c
new file mode 100644
index 00000000000..4dd8a3a7e17
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c
@@ -0,0 +1,44 @@
+#include <assert.h>
+#include <stdlib.h>
+
+struct str1 {
+  int a;
+  int b;
+  int *c;
+};
+
+#define N 1024
+
+int
+main (int argc, char *argv[])
+{
+  struct str1 s;
+
+  s.a = 1;
+  s.b = 2;
+  s.c = (int *) malloc (sizeof (int) * N);
+
+  for (int i = 0; i < N; i++)
+    s.c[i] = i + 10;
+
+  #pragma acc enter data copyin(s.a, s.b, s.c[0:N])
+
+  #pragma acc serial present(s.a, s.b, s.c[0:N]) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
+  {
+    s.a = 3;
+    s.b = 4;
+    for (int i = 0; i < N; i++)
+      s.c[i] = i + 20;
+  }
+
+  #pragma acc exit data copyout(s.a, s.b, s.c[0:N])
+
+  assert (s.a == 3);
+  assert (s.b == 4);
+  for (int i = 0; i < N; i++)
+    assert (s.c[i] == i + 20);
+
+  free (s.c);
+
+  return 0;
+}
-- 
2.23.0

[PATCH 2/3] Strip GOMP_MAP_STRUCT from OpenACC exit data mappings

[Julian Brown]

As flagged by Thomas, the GOMP_MAP_STRUCT mapping is not itself necessary
for OpenACC "exit data" directives, and is skipped over (now) in libgomp.
We might as well not emit it to start with, in line with the equivalent
OpenMP directive. We can keep the "no-op" handling in libgomp for the
reason of backward compatibility.

I've added a new scan test for the new behaviour. OK?

Julian

	gcc/
	* gimplify.c (gimplify_adjust_omp_clauses): Remove GOMP_MAP_STRUCT
	mapping from OpenACC exit data directives.
---
 gcc/gimplify.c                                |  3 +-
 .../goacc/struct-enter-exit-data-1.c          | 35 +++++++++++++++++++
 2 files changed, 37 insertions(+), 1 deletion(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index cb08b26dc65..e14932fafaf 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10396,7 +10396,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 		}
 	    }
 	  else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
-		   && code == OMP_TARGET_EXIT_DATA)
+		   && (code == OMP_TARGET_EXIT_DATA
+		       || code == OACC_EXIT_DATA))
 	    remove = true;
 	  else if (DECL_SIZE (decl)
 		   && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
diff --git a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
new file mode 100644
index 00000000000..4be5674b23e
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
@@ -0,0 +1,35 @@
+/* Check that extraneous GOMP_MAP_STRUCTs are removed from OpenACC exit data
+   directives.  */
+
+/* { dg-additional-options "-fdump-tree-omplower" } */
+
+#include <stdlib.h>
+
+struct str {
+  int a;
+  int *b;
+  int *c;
+  int d;
+};
+
+#define N 1024
+
+int
+main (int argc, char *argv[])
+{
+  struct str s;
+
+  s.b = (int *) malloc (sizeof (int) * N);
+  s.c = (int *) malloc (sizeof (int) * N);
+
+  #pragma acc enter data copyin(s.a, s.b[0:N], s.c[0:N], s.d)
+  /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.d \[len: [0-9]+\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)} omplower } } */
+
+  #pragma acc exit data copyout(s.a, s.b[0:N], s.c[0:N], s.d)
+  /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.d \[len: [0-9]+\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)} omplower } } */
+
+  free (s.b);
+  free (s.c);
+
+  return 0;
+}
-- 
2.23.0

[PATCH 3/3] Fortran derived-type mapping fix

[Julian Brown]

This patch provides the same fix for Fortran derived type component
mappings with derived types as the patch posted previously:

  https://gcc.gnu.org/legacy-ml/gcc-patches/2020-01/msg00512.html

IIUC this part was previously approved by Tobias.  OK?

Julian

	gcc/fortran/
	* trans-openmp.c (gfc_trans_omp_clauses): Use 'inner' not 'decl' for
	derived type members which themselves have derived types.

	gcc/testsuite/
	* gfortran.dg/goacc/mapping-tests-3.f90: New test.
	* gfortran.dg/goacc/mapping-tests-4.f90: New test.
---
 gcc/fortran/trans-openmp.c                      |  4 ++--
 .../gfortran.dg/goacc/mapping-tests-3.f90       | 15 +++++++++++++++
 .../gfortran.dg/goacc/mapping-tests-4.f90       | 17 +++++++++++++++++
 3 files changed, 34 insertions(+), 2 deletions(-)
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 7e2f6256c43..02c40fdc660 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2774,9 +2774,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			}
 		      else
 			{
-			  OMP_CLAUSE_DECL (node) = decl;
+			  OMP_CLAUSE_DECL (node) = inner;
 			  OMP_CLAUSE_SIZE (node)
-			    = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+			    = TYPE_SIZE_UNIT (TREE_TYPE (inner));
 			}
 		    }
 		  else if (lastcomp->next
diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
new file mode 100644
index 00000000000..312f596461e
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
@@ -0,0 +1,15 @@
+! { dg-options "-fopenacc -fdump-tree-omplower" }
+
+subroutine foo
+  type one
+    integer i, j
+  end type
+  type two
+    type(one) A, B
+  end type
+
+  type(two) x
+
+  !$acc enter data copyin(x%A)
+! { dg-final { scan-tree-dump-times "omp target oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\)" 1 "omplower" } }
+end
diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
new file mode 100644
index 00000000000..6257af942df
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
@@ -0,0 +1,17 @@
+subroutine foo
+  type one
+    integer i, j
+  end type
+  type two
+    type(one) A, B
+  end type
+
+  type(two) x
+
+! This is accepted at present, although it represents a probably-unintentional
+! overlapping subcopy.
+  !$acc enter data copyin(x%A, x%A%i)
+! But this raises an error.
+  !$acc enter data copyin(x%A, x%A%i, x%A%i)
+! { dg-error ".x.a.i. appears more than once in map clauses" "" { target "*-*-*" } 15 }
+end
-- 
2.23.0

Re: [PATCH] Fix component mappings with derived types for OpenACC

[Julian Brown]

On Tue, 19 May 2020 14:23:34 +0200
Thomas Schwinge <thomas@codesourcery.com> wrote:

> Hi!
> 
> On 2020-01-28T13:41:00+0000, Julian Brown <julian@codesourcery.com>
> wrote:
> > On Fri, 24 Jan 2020 10:58:49 +0100
> > Tobias Burnus <tobias@codesourcery.com> wrote:  
> >> the gfortran part is rather obvious and it and the test case look
> >> fine to me → OK.
> >> The oacc-mem.c also looks okay, but I assume Thomas needs to 
> >> rubber-stamp it.  
> >
> > I understand that Thomas is unavailable for the time being, so
> > won't be able to use his rubber-stamp powers. I added the offending
> > libgomp code to start with though, so I think I can go ahead and
> > commit the patch. I'll hold off for 24 hours though in case there
> > are any objections (Jakub?).  
> 
> So, in the end you didn't commit this, and now we've got the
> "un-fixed" OpenACC/Fortran code generation in the GCC 10.1 release,
> so have to deal with it in one way or another going forward,
> regarding libgomp ABI compatibility.  (Ideally), we need to make sure
> that "un-fixed" GCC 10.1-built executables continue to work "as good
> as before" when dynamically linked/running against "fixed" GCC 10.1+
> shared libraries.
> 
> Do we get such desired behavior by the patch quoted below?  In
> particular: removing the 'GOMP_MAP_STRUCT' handling code, but leaving
> in the empty 'case GOMP_MAP_STRUCT:' as a no-op, so that we don't run
> into 'default:' case 'goacc_exit_data_internal UNHANDLED kind'?  Is
> that sufficient?
> 
> Is my understanding correct that "fixed" GCC won't generate such
> 'GOMP_MAP_STRUCT' anymore (I have't studied in detail), and this empty
> 'case GOMP_MAP_STRUCT:' only remains in here for backwards
> compatibility? In this case, please add a comment to the code,
> stating this.  Otherwise, please add a comment why "do nothing" is
> appropriate for 'GOMP_MAP_STRUCT'.  In particular, for both
> scenarios, why we don't need to skip the following 'sizes[i]'
> mappings?

I'm not sure if I got the threading right, but I've now followed up on
this discussion here:

  https://gcc.gnu.org/pipermail/gcc-patches/2020-June/547300.html

Thanks,

Julian

Add 'libgomp.oacc-c-c++-common/struct-copyout-{1, 2}.c' (was: [PATCH 1/3] OpenACC "exit data" copyout for struct members)

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 879 bytes --]

Hi!

On 2020-06-04T06:40:53-0700, Julian Brown <julian@codesourcery.com> wrote:
> [...] I've added a couple of new testcases to verify the behaviour
> (which fail without the patch).

Thanks.  Given the preparational patches pushed yesterday, these now PASS
already.

>       * testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c: New test.
>       * testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c: New test.

Pushed "Add 'libgomp.oacc-c-c++-common/struct-copyout-{1,2}.c'" to master
branch in commit 9643f5bbe237764cbefc975e934d1281f47ee3c2, and
releases/gcc-10 branch in commit
52d737058897eb438099b57234d41330147d0b6f, see attached.


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Add-libgomp.oacc-c-c-common-struct-copyout-1-2-.c.patch --]
[-- Type: text/x-diff, Size: 2753 bytes --]

From 9643f5bbe237764cbefc975e934d1281f47ee3c2 Mon Sep 17 00:00:00 2001
From: Julian Brown <julian@codesourcery.com>
Date: Thu, 4 Jun 2020 06:40:53 -0700
Subject: [PATCH] Add 'libgomp.oacc-c-c++-common/struct-copyout-{1,2}.c'

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c: New test.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
---
 .../struct-copyout-1.c                        | 38 ++++++++++++++++
 .../struct-copyout-2.c                        | 44 +++++++++++++++++++
 2 files changed, 82 insertions(+)
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
new file mode 100644
index 000000000000..b86f1c921a98
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
@@ -0,0 +1,38 @@
+#include <assert.h>
+
+struct str1 {
+  int a;
+  int b;
+};
+
+struct str2 {
+  int c;
+  int d;
+  struct str1 s;
+};
+
+int
+main (int argc, char *argv[])
+{
+  struct str2 t;
+
+  t.c = 1;
+  t.d = 2;
+  t.s.a = 3;
+  t.s.b = 4;
+
+  #pragma acc enter data copyin(t.s)
+
+  #pragma acc serial present(t.s) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
+  {
+    t.s.a = 5;
+    t.s.b = 6;
+  }
+
+  #pragma acc exit data copyout(t.s)
+
+  assert (t.s.a == 5);
+  assert (t.s.b == 6);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c
new file mode 100644
index 000000000000..4dd8a3a7e175
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c
@@ -0,0 +1,44 @@
+#include <assert.h>
+#include <stdlib.h>
+
+struct str1 {
+  int a;
+  int b;
+  int *c;
+};
+
+#define N 1024
+
+int
+main (int argc, char *argv[])
+{
+  struct str1 s;
+
+  s.a = 1;
+  s.b = 2;
+  s.c = (int *) malloc (sizeof (int) * N);
+
+  for (int i = 0; i < N; i++)
+    s.c[i] = i + 10;
+
+  #pragma acc enter data copyin(s.a, s.b, s.c[0:N])
+
+  #pragma acc serial present(s.a, s.b, s.c[0:N]) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
+  {
+    s.a = 3;
+    s.b = 4;
+    for (int i = 0; i < N; i++)
+      s.c[i] = i + 20;
+  }
+
+  #pragma acc exit data copyout(s.a, s.b, s.c[0:N])
+
+  assert (s.a == 3);
+  assert (s.b == 4);
+  for (int i = 0; i < N; i++)
+    assert (s.c[i] == i + 20);
+
+  free (s.c);
+
+  return 0;
+}
-- 
2.26.2


[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #3: 0001-Add-libgomp.oacc-c-c-common-struct-copyout-1-2-..g10.patch --]
[-- Type: text/x-diff, Size: 2822 bytes --]

From 52d737058897eb438099b57234d41330147d0b6f Mon Sep 17 00:00:00 2001
From: Julian Brown <julian@codesourcery.com>
Date: Thu, 4 Jun 2020 06:40:53 -0700
Subject: [PATCH] Add 'libgomp.oacc-c-c++-common/struct-copyout-{1,2}.c'

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c: New test.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
(cherry picked from commit 9643f5bbe237764cbefc975e934d1281f47ee3c2)
---
 .../struct-copyout-1.c                        | 38 ++++++++++++++++
 .../struct-copyout-2.c                        | 44 +++++++++++++++++++
 2 files changed, 82 insertions(+)
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
new file mode 100644
index 000000000000..b86f1c921a98
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
@@ -0,0 +1,38 @@
+#include <assert.h>
+
+struct str1 {
+  int a;
+  int b;
+};
+
+struct str2 {
+  int c;
+  int d;
+  struct str1 s;
+};
+
+int
+main (int argc, char *argv[])
+{
+  struct str2 t;
+
+  t.c = 1;
+  t.d = 2;
+  t.s.a = 3;
+  t.s.b = 4;
+
+  #pragma acc enter data copyin(t.s)
+
+  #pragma acc serial present(t.s) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
+  {
+    t.s.a = 5;
+    t.s.b = 6;
+  }
+
+  #pragma acc exit data copyout(t.s)
+
+  assert (t.s.a == 5);
+  assert (t.s.b == 6);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c
new file mode 100644
index 000000000000..4dd8a3a7e175
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c
@@ -0,0 +1,44 @@
+#include <assert.h>
+#include <stdlib.h>
+
+struct str1 {
+  int a;
+  int b;
+  int *c;
+};
+
+#define N 1024
+
+int
+main (int argc, char *argv[])
+{
+  struct str1 s;
+
+  s.a = 1;
+  s.b = 2;
+  s.c = (int *) malloc (sizeof (int) * N);
+
+  for (int i = 0; i < N; i++)
+    s.c[i] = i + 10;
+
+  #pragma acc enter data copyin(s.a, s.b, s.c[0:N])
+
+  #pragma acc serial present(s.a, s.b, s.c[0:N]) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
+  {
+    s.a = 3;
+    s.b = 4;
+    for (int i = 0; i < N; i++)
+      s.c[i] = i + 20;
+  }
+
+  #pragma acc exit data copyout(s.a, s.b, s.c[0:N])
+
+  assert (s.a == 3);
+  assert (s.b == 4);
+  for (int i = 0; i < N; i++)
+    assert (s.c[i] == i + 20);
+
+  free (s.c);
+
+  return 0;
+}
-- 
2.26.2

Re: [PATCH 3/3] Fortran derived-type mapping fix

[Thomas Schwinge]

Hi Julian!

On 2020-06-04T06:40:55-0700, Julian Brown <julian@codesourcery.com> wrote:
> This patch provides the same fix for Fortran derived type component
> mappings with derived types as the patch posted previously:
>
>   https://gcc.gnu.org/legacy-ml/gcc-patches/2020-01/msg00512.html
>
> IIUC this part was previously approved by Tobias.  OK?

I don't have in-depth knowledge of these Fortran front end OMP parts.
Tobias had reviewed and approved your patch in January, and I'll assume
his approval still holds.  Or are you asking for further review of the
patch?

I suppose we want this backported to releases/gcc-10 branch, too?

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

As '-fopenacc' is the default for 'gfortran.dg/goacc/', you can avoid to
repeat it by using '{ dg-additional-options "-fdump-tree-omplower" }'.

Just curious: specific reason to scan '-fdump-tree-omplower' -- wouldn't
the problem be visible in '-fdump-tree-original' or
'-fdump-tree-gimplify' already?

> +
> +subroutine foo
> +  type one
> +    integer i, j
> +  end type
> +  type two
> +    type(one) A, B
> +  end type
> +
> +  type(two) x
> +
> +  !$acc enter data copyin(x%A)
> +! { dg-final { scan-tree-dump-times "omp target oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\)" 1 "omplower" } }
> +end

> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90
> @@ -0,0 +1,17 @@
> +subroutine foo
> +  type one
> +    integer i, j
> +  end type
> +  type two
> +    type(one) A, B
> +  end type
> +
> +  type(two) x
> +
> +! This is accepted at present, although it represents a probably-unintentional
> +! overlapping subcopy.
> +  !$acc enter data copyin(x%A, x%A%i)

Should that missing diagnostic get a PR filed, and that one be referenced
here?

> +! But this raises an error.
> +  !$acc enter data copyin(x%A, x%A%i, x%A%i)
> +! { dg-error ".x.a.i. appears more than once in map clauses" "" { target "*-*-*" } 15 }
> +end

Not a problem, but unusual to see '"*-*-*"' quoted.  Preferable to use
relative '.-1' instead of absolute '15'.


Grüße
 Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter