Hi Julian, Tobias!
On 2020-07-27T15:33:41+0100, Julian Brown <jul...@codesourcery.com> wrote:
> On Fri, 17 Jul 2020 13:16:11 +0200
> Thomas Schwinge <tho...@codesourcery.com> wrote:
>> 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.
ACK. Do you remember if you already had a testcase (conceptual, or
actual) to demonstrate that problem?
>> >>> 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.
(Just remains the question who's going to do that, "later"...) ;-\
>> >> 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.
ACK -- good to hear that this is the actual solution here.
>> 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.
I'm attaching an incremental patch (I have tested that) to merge the
testcases into one file, and make it work on current master branch
without the pending "Refuse update/copyout for blocks with attached
pointers" changes. (We then later have to adjust the testcase here as
part of these changes.)
> Tested with offloading to NVPTX. OK?
Thanks. OK for master and releases/gcc-10 branches from my point of
view, but maybe Tobias can also have a look, please; two
comments/suggestions:
> 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".
(Need to remove that last sentence.)
> --- 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. */
I don't understand this code good enough to be sure that 'goto
finalize_map_clause' doesn't skip anything we may actually need -- for
the many "special" cases that Fortran has. Is it the case that it's the
correct thing to do, given that we're skipping 'node' completely.
I just had an idea how to make that clearer (maybe?) (untested, of
course): instead of the 'node', 'node2', 'node3' shuffling and 'goto
finalize_map_clause', don't do the shuffling and instead 'goto
finalize_map_clause_auxilliary' (better name maybe?):
finalize_map_clause:
omp_clauses = gfc_trans_add_clause (node, omp_clauses);
+finalize_map_clause_auxilliary:
if (node2)
omp_clauses = gfc_trans_add_clause (node2, omp_clauses);
if (node3)
omp_clauses = gfc_trans_add_clause (node3, omp_clauses);
if (node4)
omp_clauses = gfc_trans_add_clause (node4, omp_clauses);
(Just an idea; can also be done separately, later.)
> --- 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. */
Maybe be good to add a comment why it's OK to do nothing for
'GOMP_MAP_TO_PSET'?
Grüße
Thomas
> --- 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 "" }
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander
Walter
>From e3241486f68c077006513ea41c59ba3fdaeca7f7 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Wed, 29 Jul 2020 15:57:17 +0200
Subject: [PATCH] into: openacc: No attach/detach present/release mappings for
array descriptors
---
.../attach-descriptor-1.f90 | 93 ++++++++++++++++---
.../attach-descriptor-3.f90 | 68 --------------
.../attach-descriptor-4.f90 | 61 ------------
3 files changed, 80 insertions(+), 142 deletions(-)
delete mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90
delete mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
index 9f159fa3b75..960b9f94507 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
@@ -1,9 +1,10 @@
! { dg-do run }
! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
-program att
+subroutine test(variant)
use openacc
implicit none
+ integer :: variant
type t
integer :: arr1(10)
integer, allocatable :: arr2(:)
@@ -27,31 +28,97 @@ program att
myptr => tarr
- !$acc enter data attach(myvar%arr2, myptr)
+ if (variant == 0 &
+ .or. variant == 3 &
+ .or. variant == 5) then
+ !$acc enter data attach(myvar%arr2, myptr)
+ else if (variant == 1 &
+ .or. variant == 2 &
+ .or. variant == 4) then
+ !$acc enter data attach(myvar%arr2, myptr)
+ !$acc enter data attach(myvar%arr2, myptr)
+ else
+ ! Internal error.
+ stop 1
+ end if
! FIXME: This warning is emitted on the wrong line number.
- ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 39 }
+ ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 52 }
!$acc serial present(myvar%arr2)
do i=1,10
- myvar%arr1(i) = i
- myvar%arr2(i) = i
+ myvar%arr1(i) = i + variant
+ myvar%arr2(i) = i - variant
end do
- myptr(3) = 99
+ myptr(3) = 99 - variant
!$acc end serial
- !$acc exit data detach(myvar%arr2, myptr)
+ if (variant == 0) then
+ !$acc exit data detach(myvar%arr2, myptr)
+ else if (variant == 1) then
+ !$acc exit data detach(myvar%arr2, myptr)
+ !$acc exit data detach(myvar%arr2, myptr)
+ else if (variant == 2) then
+ !$acc exit data detach(myvar%arr2, myptr)
+ !$acc exit data detach(myvar%arr2, myptr) finalize
+ else if (variant == 3 &
+ .or. variant == 4) then
+ !$acc exit data detach(myvar%arr2, myptr) finalize
+ else if (variant == 5) then
+ ! Do not detach.
+ else
+ ! Internal error.
+ stop 2
+ end if
+
+ 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 10
+ 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)) stop 11
+ 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(tarr)) stop 12
+ 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
+ if (myvar%arr1(i) .ne. i + variant) stop 50
+ if (variant == 5) then
+ ! We have not detached, so have copyied out a device pointer, so cannot
+ ! access 'myvar%arr2' on the host.
+ else
+ if (myvar%arr2(i) .ne. i - variant) stop 51
+ end if
end do
- if (tarr(3) .ne. 99) stop 3
+ if (tarr(3) .ne. 99 - variant) stop 52
+
+ if (variant == 5) then
+ ! If not explicitly stopping here, we'd in the following try to deallocate
+ ! the device pointer on the host, SIGSEGV.
+ stop
+ end if
+end subroutine test
+
+program att
+ implicit none
+
+ call test(0)
+
+ call test(1)
+
+ call test(2)
+
+ call test(3)
+
+ call test(4)
+ call test(5)
+ ! Make sure that 'test(5)' has stopped the program.
+ stop 60
end program att
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90
deleted file mode 100644
index f0e57b47453..00000000000
--- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90
+++ /dev/null
@@ -1,68 +0,0 @@
-! { 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
deleted file mode 100644
index 9dbf53d0213..00000000000
--- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90
+++ /dev/null
@@ -1,61 +0,0 @@
-! { 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.17.1
