On Fri, 17 Jul 2020 13:16:11 +0200
Thomas Schwinge <tho...@codesourcery.com> wrote:
> Hi Julian, Tobias!
>
> On 2020-07-15T12:28:42+0200, Thomas Schwinge
> <tho...@codesourcery.com> wrote:
> > On 2020-07-14T13:43:37+0200, I wrote:
> >> On 2020-06-16T15:39:44-0700, Julian Brown
> >> <jul...@codesourcery.com> wrote:
> >>> As mentioned in the blurb for the previous patch, an "attach"
> >>> operation for a Fortran pointer with an array descriptor must
> >>> copy that array descriptor to the target.
> >>
> >> Heh, I see -- I don't think I had read the OpenACC standard in
> >> that way, but I think I agree your interpretation is fine.
> >>
> >> This does not create some sort of memory leak -- everything
> >> implicitly allocated there will eventually be deallocated again,
> >> right?
>
> Unanswered -- but I may now have found this problem, and also found
> "the reverse problem" ('finalize'); see below.
Sorry, I didn't answer this explicitly -- the idea was to pair alloc
(present) and release mappings for the pointed-to data. In that way,
the idea was for the release mapping to perform that deallocation. That
was partly done so that the existing handling in gfc_trans_omp_clauses
could be used for this case without too much disruption to the code --
but actually, after Tobias's reorganisation of that function, that's
not really so much of an issue any more.
You can still get a "leak" if you try to attach a synthesized/temporary
array descriptor that goes out of scope before the pointed-to data it
refers to does -- that's a problem I've mentioned earlier, and is
kind-of unavoidable unless we do some more sophisticated analysis to
diagnose it as user error.
> >>> This patch arranges for that to be so.
> >>
> >> In response to the new OpenACC/Fortran testcase that I'd submtited
> >> in
> >> <87wo3co0tm.fsf@euler.schwinge.homeip.net">http://mid.mail-archive.com/87wo3co0tm.fsf@euler.schwinge.homeip.net>,
> >> you (Julian) correctly supposed in
> >> <http://mid.mail-archive.com/20200709223246.23a4d0e0@squid.athome>,
> >> that this patch indeed does resolve that testcase, too. That
> >> wasn't obvious to me. So, similar to
> >> 'libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-{1.2}.c',
> >> please include my new OpenACC/Fortran testcase (if that makes
> >> sense to you), and reference PR95270 in the commit log.
> >
> > My new OpenACC/Fortran testcase got again broken ('libgomp: pointer
> > target not mapped for attach') by Tobias' commit
> > 102502e32ea4e8a75d6b252ba319d09d735d9aa7 "[OpenMP, Fortran] Add
> > structure/derived-type element mapping",
> > <c5b43e02-d1d5-e7cf-c11c-6daf1e8f33c5@codesourcery.com">http://mid.mail-archive.com/c5b43e02-d1d5-e7cf-c11c-6daf1e8f33c5@codesourcery.com>.
> >
> > Similar ('libgomp: attempt to attach null pointer') for your new
> > 'libgomp.oacc-fortran/attach-descriptor-1.f90'.
> >
> > (Whether or not 'attach'ing 'NULL' should actually be allowed, is a
> > separate topic for discussion.)
> >
> > So this patch here will (obviously) need to be adapted to what
> > Tobias changed.
>
> I see what you pushed in commit
> 39dda0020801045d9a604575b2a2593c05310015 "openacc: Fix standalone
> attach for Fortran assumed-shape array pointers" indeed has become
> much smaller/simpler. :-)
Yes, thank you.
> (But, (parts of?) Tobias' commit mentioned above (plus commit
> 524862db444b6544c6dc87c5f06f351100ecf50d "Fix goacc/finalize-1.f tree
> dump-scanning for -m32", if applicable) will then also need to be
> backported to releases/gcc-10 branch (once un-frozen).)
>
> > (Plus my more general questions quoted above and below.)
>
> >>> OK?
> >>
> >> Basically yes (for master and releases/gcc-10 branches), but please
> >> consider the following:
> >>
> >>> --- a/gcc/fortran/trans-openmp.c
> >>> +++ b/gcc/fortran/trans-openmp.c
> >>> @@ -2573,8 +2573,44 @@ gfc_trans_omp_clauses (stmtblock_t *block,
> >>> gfc_omp_clauses *clauses, }
> >>> }
> >>> if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))
> >>> - && n->u.map_op != OMP_MAP_ATTACH
> >>> - && n->u.map_op != OMP_MAP_DETACH)
> >>> + && (n->u.map_op == OMP_MAP_ATTACH
> >>> + || n->u.map_op == OMP_MAP_DETACH))
> >>> + {
> >>> + tree type = TREE_TYPE (decl);
> >>> + tree data = gfc_conv_descriptor_data_get
> >>> (decl);
> >>> + if (present)
> >>> + data = gfc_build_cond_assign_expr
> >>> (block, present,
> >>> + data,
> >>> +
> >>> null_pointer_node);
> >>> + tree ptr
> >>> + = fold_convert (build_pointer_type
> >>> (char_type_node),
> >>> + data);
> >>> + ptr = build_fold_indirect_ref (ptr);
> >>> + /* Standalone attach clauses used with
> >>> arrays with
> >>> + descriptors must copy the descriptor to
> >>> the target,
> >>> + else they won't have anything to
> >>> perform the
> >>> + attachment onto (see OpenACC 2.6,
> >>> "2.6.3. Data
> >>> + Structures with Pointers"). */
> >>> + OMP_CLAUSE_DECL (node) = ptr;
> >>> + node2 = build_omp_clause (input_location,
> >>> OMP_CLAUSE_MAP);
> >>> + OMP_CLAUSE_SET_MAP_KIND (node2,
> >>> GOMP_MAP_TO_PSET);
> >>> + OMP_CLAUSE_DECL (node2) = decl;
> >>> + OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT
> >>> (type);
> >>> + node3 = build_omp_clause (input_location,
> >>> OMP_CLAUSE_MAP);
> >>> + if (n->u.map_op == OMP_MAP_ATTACH)
> >>> + {
> >>> + OMP_CLAUSE_SET_MAP_KIND (node3,
> >>> GOMP_MAP_ATTACH);
> >>> + n->u.map_op = OMP_MAP_ALLOC;
> >>> + }
> >>> + else /* OMP_MAP_DETACH. */
> >>> + {
> >>> + OMP_CLAUSE_SET_MAP_KIND (node3,
> >>> GOMP_MAP_DETACH);
> >>> + n->u.map_op = OMP_MAP_RELEASE;
> >>> + }
> >>> + OMP_CLAUSE_DECL (node3) = data;
> >>> + OMP_CLAUSE_SIZE (node3) = size_int (0);
> >>> + }
> >>
> >> So this ("case A") duplicates most of the code from...
> >>
> >>> + else if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE
> >>> (decl))) {
> >>> [...]
> >>
> >> ... this existing case here ("case B"). It's not clear to me if
> >> these two cases really still need to be handled separately, and a
> >> little bit differently (regarding 'if (present)' handling, for
> >> example), or if they could/should (?) be merged? Tobias, do you
> >> have an opinion?
>
> (These have been merged.)
>
> >> Do we have sufficient testsuite coverage? (For example,
> >> 'attach'/'detach' with 'present == false', if that makes sense, or
> >> any other thing that case A is doing differently from case B?)
>
> (I'm not sure we're actually testing all relevant cases.)
...probably still not, sorry... more tests can be added later though of
course.
> >> Shouldn't
> >> this get '-fdump-tree-original' and/or '-fdump-tree-gimple'
> >> testcases, similar to 'gfortran.dg/goacc/finalize-1.f', so that we
> >> verify/document what we generate here?
>
> So I guess I had -- unconsciously? ;-) -- mentioned
> -fdump-tree-gimple' and 'gfortran.dg/goacc/finalize-1.f' for a
> reason. That displays how the 'finalize' clause is implemented (see
> WIP patch attached, 'gfortran.dg/goacc/attach-descriptor.f90'), and...
[snip]
> What should happen in this case? Do we agree that 'exit data
> detach(myptr)' should *never* unmap 'myptr => tarr', but really should
> just unmap the 'myptr' array descriptor?
>
> We can add special handling so that for standalone 'detach', a
> 'finalize' doesn't turn 'release' into 'delete', but that doesn't
> feel like the correct solution.
I don't think we actually need the alloc/release (with the latter turned
into "delete" for finalize) at all -- we just need to map the array
descriptor and perform the attach (or detach) as necessary. That's what
the attached patch does. Then, the pointed-to data's reference counts,
etc. will not be modified by attach/detach operations at all.
> Also, we have a different -- bigger? -- problem: if we, for example,
> 'attach(myptr)' twice, that operation will include twice times
> incrementing the reference count of 'myptr => tarr', and that'll then
> conflict with a 'copyout(myptr)', as that one then sees unexpected
> reference counts. That's a different variant of the "[OpenACC] Deep
> copy attach/detach should not affect reference counts" problem?
>
> Basically (see WIP patch attached,
> 'libgomp.oacc-fortran/attach-descriptor-1_.f90'):
Hmm, yes -- FWIW, this is caught by the "Refuse update/copyout for
blocks with attached pointers" patch. (In fact the attached patch
assumes that patch is already committed -- else the
attach-descriptor-4.f90 test should be XFAILed or omitted). So if we
want that one, this problem is sidestepped, I think.
Tested with offloading to NVPTX. OK?
Thanks,
Julian
>From d53e4f1cd450062163e7e96a469c2f56cfac65ee Mon Sep 17 00:00:00 2001
From: Julian Brown <jul...@codesourcery.com>
Date: Mon, 27 Jul 2020 06:29:02 -0700
Subject: [PATCH] openacc: No attach/detach present/release mappings for array
descriptors
Standalone attach and detach clauses should not create present/release
mappings for Fortran array descriptors (e.g. used when we have a pointer
to an array), both because it is unnecessary and because those mappings
will be incorrectly subject to reference counting. Simply omitting the
mappings means we just use GOMP_MAP_TO_PSET and GOMP_MAP_{ATTACH,DETACH}
mappings for array descriptors.
That requires a tweak in gimplify.c, since we may now see GOMP_MAP_TO_PSET
without a preceding data-movement mapping.
The new attach-descriptor-4.f90 test relies on the checking performed
by the patch "Refuse update/copyout for blocks with attached pointers".
2020-07-27 Julian Brown <jul...@codesourcery.com>
Thomas Schwinge <tho...@codesourcery.com>
gcc/fortran/
* trans-openmp.c (gfc_trans_omp_clauses): Don't create present/release
mappings for array descriptors.
gcc/
* gimplify.c (gimplify_omp_target_update): Allow GOMP_MAP_TO_PSET
without a preceding data-movement mapping.
gcc/testsuite/
* gfortran.dg/goacc/attach-descriptor.f90: Update pattern output. Add
scanning of gimplify dump.
libgomp/
* testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90: Don't run for
shared-memory devices. Add more checking.
* testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90: New test.
* testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90: New test.
Co-Authored-By: Thomas Schwinge <tho...@codesourcery.com>
---
gcc/fortran/trans-openmp.c | 44 +++++++-----
gcc/gimplify.c | 3 +-
.../gfortran.dg/goacc/attach-descriptor.f90 | 17 ++++-
.../attach-descriptor-1.f90 | 6 +-
.../attach-descriptor-3.f90 | 68 +++++++++++++++++++
.../attach-descriptor-4.f90 | 61 +++++++++++++++++
6 files changed, 177 insertions(+), 22 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index d12d7fbddac..1a8f3277de3 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2718,23 +2718,6 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
- if (n->u.map_op == OMP_MAP_ATTACH)
- {
- /* Standalone attach clauses used with arrays with
- descriptors must copy the descriptor to the target,
- else they won't have anything to perform the
- attachment onto (see OpenACC 2.6, "2.6.3. Data
- Structures with Pointers"). */
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
- }
- else if (n->u.map_op == OMP_MAP_DETACH)
- {
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
- }
- else
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
if (present)
{
ptr = gfc_conv_descriptor_data_get (decl);
@@ -2748,6 +2731,33 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
OMP_CLAUSE_DECL (node3)
= gfc_conv_descriptor_data_get (decl);
OMP_CLAUSE_SIZE (node3) = size_int (0);
+ if (n->u.map_op == OMP_MAP_ATTACH)
+ {
+ /* Standalone attach clauses used with arrays with
+ descriptors must copy the descriptor to the target,
+ else they won't have anything to perform the
+ attachment onto (see OpenACC 2.6, "2.6.3. Data
+ Structures with Pointers"). */
+ OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
+ /* We don't want to map PTR at all in this case, so
+ delete its node and shuffle the others down. */
+ node = node2;
+ node2 = node3;
+ node3 = NULL;
+ goto finalize_map_clause;
+ }
+ else if (n->u.map_op == OMP_MAP_DETACH)
+ {
+ OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
+ /* Similarly to above, we don't want to unmap PTR
+ here. */
+ node = node2;
+ node2 = node3;
+ node3 = NULL;
+ goto finalize_map_clause;
+ }
+ else
+ OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
/* We have to check for n->sym->attr.dimension because
of scalar coarrays. */
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 15dfee903ab..f4c31d2870d 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -13013,8 +13013,9 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
have_clause = true;
break;
- case GOMP_MAP_POINTER:
case GOMP_MAP_TO_PSET:
+ break;
+ case GOMP_MAP_POINTER:
/* TODO PR92929: we may see these here, but they'll always follow
one of the clauses above, and will be handled by libgomp as
one group, so no handling required here. */
diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
index 9ca36f770c7..373bdcb2114 100644
--- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
@@ -1,4 +1,4 @@
-! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
program att
implicit none
@@ -11,8 +11,19 @@ program att
integer, pointer :: myptr(:)
!$acc enter data attach(myvar%arr2, myptr)
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(alloc:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
!$acc exit data detach(myvar%arr2, myptr)
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(release:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
+
+ ! Test valid usage and processing of the finalize clause.
+ !$acc exit data detach(myvar%arr2, myptr) finalize
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize;$" 1 "original" } }
+ ! For array-descriptor detaches, we no longer generate a "release" mapping
+ ! for the pointed-to data for gimplify.c to turn into "delete". Make sure
+ ! the mapping still isn't there.
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
+
end program att
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
index 5d79cbc14fc..9f159fa3b75 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
@@ -1,4 +1,5 @@
! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
program att
use openacc
@@ -29,7 +30,7 @@ program att
!$acc enter data attach(myvar%arr2, myptr)
! FIXME: This warning is emitted on the wrong line number.
- ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 38 }
+ ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 39 }
!$acc serial present(myvar%arr2)
do i=1,10
myvar%arr1(i) = i
@@ -41,8 +42,11 @@ program att
!$acc exit data detach(myvar%arr2, myptr)
call acc_copyout(myvar%arr2)
+ if (acc_is_present(myvar%arr2)) stop 10
call acc_copyout(myvar)
+ if (acc_is_present(myvar)) stop 11
call acc_copyout(tarr)
+ if (acc_is_present(tarr)) stop 12
do i=1,10
if (myvar%arr1(i) .ne. i) stop 1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90
new file mode 100644
index 00000000000..f0e57b47453
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90
@@ -0,0 +1,68 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+program att
+ use openacc
+ implicit none
+ type t
+ integer :: arr1(10)
+ integer, allocatable :: arr2(:)
+ end type t
+ integer :: i
+ type(t) :: myvar
+ integer, target :: tarr(10)
+ integer, pointer :: myptr(:)
+
+ allocate(myvar%arr2(10))
+
+ do i=1,10
+ myvar%arr1(i) = 0
+ myvar%arr2(i) = 0
+ tarr(i) = 0
+ end do
+
+ call acc_copyin(myvar)
+ call acc_copyin(myvar%arr2)
+ call acc_copyin(tarr)
+
+ myptr => tarr
+
+ !$acc enter data attach(myvar%arr2, myptr)
+ !$acc enter data attach(myvar%arr2, myptr)
+
+ ! FIXME: This warning is emitted on the wrong line number.
+ ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 40 }
+ !$acc serial present(myvar%arr2)
+ do i=1,10
+ myvar%arr1(i) = i
+ myvar%arr2(i) = i
+ end do
+ myptr(3) = 99
+ !$acc end serial
+
+ !$acc exit data detach(myvar%arr2, myptr) finalize
+
+ if (.not. acc_is_present(myvar%arr2)) stop 10
+ if (.not. acc_is_present(myvar)) stop 11
+ if (.not. acc_is_present(tarr)) stop 12
+
+ call acc_copyout(myvar%arr2)
+ if (acc_is_present(myvar%arr2)) stop 20
+ if (.not. acc_is_present(myvar)) stop 21
+ if (.not. acc_is_present(tarr)) stop 22
+ call acc_copyout(myvar)
+ if (acc_is_present(myvar%arr2)) stop 30
+ if (acc_is_present(myvar)) stop 31
+ if (.not. acc_is_present(tarr)) stop 32
+ call acc_copyout(tarr)
+ if (acc_is_present(myvar%arr2)) stop 40
+ if (acc_is_present(myvar)) stop 41
+ if (acc_is_present(tarr)) stop 42
+
+ do i=1,10
+ if (myvar%arr1(i) .ne. i) stop 1
+ if (myvar%arr2(i) .ne. i) stop 2
+ end do
+ if (tarr(3) .ne. 99) stop 3
+
+end program att
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90
new file mode 100644
index 00000000000..9dbf53d0213
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90
@@ -0,0 +1,61 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+program att
+ use openacc
+ implicit none
+ type t
+ integer :: arr1(10)
+ integer, allocatable :: arr2(:)
+ end type t
+ integer :: i
+ type(t) :: myvar
+ integer, target :: tarr(10)
+ integer, pointer :: myptr(:)
+
+ allocate(myvar%arr2(10))
+
+ do i=1,10
+ myvar%arr1(i) = 0
+ myvar%arr2(i) = 0
+ tarr(i) = 0
+ end do
+
+ call acc_copyin(myvar)
+ call acc_copyin(myvar%arr2)
+ call acc_copyin(tarr)
+
+ myptr => tarr
+
+ !$acc enter data attach(myvar%arr2, myptr)
+ !$acc enter data attach(myvar%arr2, myptr)
+
+ ! FIXME: This warning is emitted on the wrong line number.
+ ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 40 }
+ !$acc serial present(myvar%arr2)
+ do i=1,10
+ myvar%arr1(i) = i
+ myvar%arr2(i) = i
+ end do
+ myptr(3) = 99
+ !$acc end serial
+
+ !$acc exit data detach(myvar%arr2, myptr)
+
+ call acc_copyout(myvar%arr2)
+ ! { dg-output ".*copyout of block \\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\] with attached pointers(\n|\r\n|\r)+" }
+ if (acc_is_present(myvar%arr2)) stop 10
+ call acc_copyout(myvar)
+ if (acc_is_present(myvar)) stop 11
+ call acc_copyout(tarr)
+ if (acc_is_present(tarr)) stop 12
+
+ do i=1,10
+ if (myvar%arr1(i) .ne. i) stop 1
+ if (myvar%arr2(i) .ne. i) stop 2
+ end do
+ if (tarr(3) .ne. 99) stop 3
+
+end program att
+
+! { dg-shouldfail "" }
--
2.23.0
