| Message ID | 0193ff08d4a4a2b6ca86c7a891b8ff35203fa440.1592343757.git.julian@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | Refcounting and manual deep copy improvements | expand |
Hi Julian, Tobias! On 2020-06-16T15:39:44-0700, Julian Brown <julian@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? > This patch arranges for that to be so. In response to the new OpenACC/Fortran testcase that I'd submtited in <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. > 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? 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?) 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? > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 > @@ -0,0 +1,51 @@ > +program att Please add 'dg-do run', and... > + 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) > + > + ! FIXME: This warning is emitted on the wrong line number. > + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 36 } ... don't forget to adjust "36" here. ;-) > + !$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) > + call acc_copyout(myvar) > + call acc_copyout(tarr) > + > + 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 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
Hi Julian, Tobias! On 2020-07-14T13:43:37+0200, I wrote: > On 2020-06-16T15:39:44-0700, Julian Brown <julian@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? > >> This patch arranges for that to be so. > > In response to the new OpenACC/Fortran testcase that I'd submtited in > <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", <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. (Plus my more general questions quoted above and below.) Grüße Thomas >> 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? > > 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?) 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? > > >> --- /dev/null >> +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 >> @@ -0,0 +1,51 @@ >> +program att > > Please add 'dg-do run', and... > >> + 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) >> + >> + ! FIXME: This warning is emitted on the wrong line number. >> + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 36 } > > ... don't forget to adjust "36" here. ;-) > >> + !$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) >> + call acc_copyout(myvar) >> + call acc_copyout(tarr) >> + >> + 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 > > > 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
Hi Julian, Tobias! On 2020-07-15T12:28:42+0200, Thomas Schwinge <thomas@codesourcery.com> wrote: > On 2020-07-14T13:43:37+0200, I wrote: >> On 2020-06-16T15:39:44-0700, Julian Brown <julian@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. >>> This patch arranges for that to be so. >> >> In response to the new OpenACC/Fortran testcase that I'd submtited in >> <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", > <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. :-) (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.) >> 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... >>> --- /dev/null >>> +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 >>> @@ -0,0 +1,51 @@ >>> +program att >> >> Please add 'dg-do run', and... >> >>> + 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) >>> + >>> + ! FIXME: This warning is emitted on the wrong line number. >>> + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 36 } >> >> ... don't forget to adjust "36" here. ;-) >> >>> + !$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) ..., if we here 'detach' with 'finalize' added, that will turn into a 'delete' (instead of 'release') of 'myptr => tarr', and thus... >>> + >>> + call acc_copyout(myvar%arr2) >>> + call acc_copyout(myvar) >>> + call acc_copyout(tarr) >>> + >>> + 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 ..., here we won't see the updated 'tarr(3) == 99', and fail. >>> + >>> +end program att Alternativly, we can show the problem with 'acc_is_present', as in my WIP patch attached, 'libgomp.oacc-fortran/attach-descriptor-1__.f90'. (But when experimenting with 'acc_is_present' and Fortran 'pointer's, beware of PR96080 "OpenACC/Fortran runtime library routines vs. Fortran 'pointer'".) 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. 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'): call acc_copyin(tarr) ! 'rc(tarr) == 1' myptr => tarr !$acc enter data attach(myptr) ! 'rc(tarr) == 2'! (not intended by the user) !$acc enter data attach(myptr) ! 'rc(tarr) == 3'! (not intended by the user) [...] call acc_copyout(tarr) ! won't copyout, because still 'rc(tarr) = 2'! (not intended by the user) if (acc_is_present(tarr)) stop 12 ! fails Ugh. :-( Or am I confused now? 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
On Fri, 17 Jul 2020 13:16:11 +0200 Thomas Schwinge <thomas@codesourcery.com> wrote: > Hi Julian, Tobias! > > On 2020-07-15T12:28:42+0200, Thomas Schwinge > <thomas@codesourcery.com> wrote: > > On 2020-07-14T13:43:37+0200, I wrote: > >> On 2020-06-16T15:39:44-0700, Julian Brown > >> <julian@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 > >> <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", > > <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
Hi Julian, Tobias! On 2020-07-27T15:33:41+0100, Julian Brown <julian@codesourcery.com> wrote: > On Fri, 17 Jul 2020 13:16:11 +0200 > Thomas Schwinge <thomas@codesourcery.com> wrote: >> On 2020-07-15T12:28:42+0200, Thomas Schwinge >> <thomas@codesourcery.com> wrote: >> > On 2020-07-14T13:43:37+0200, I wrote: >> >> On 2020-06-16T15:39:44-0700, Julian Brown >> >> <julian@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 >> >> <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", >> > <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 <julian@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
On Thu, 30 Jul 2020 11:53:10 +0200 Thomas Schwinge <thomas@codesourcery.com> wrote: > Hi Julian, Tobias! > > On 2020-07-27T15:33:41+0100, Julian Brown <julian@codesourcery.com> > wrote: > > 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? I have the attached, but it's not "clean", i.e. not really testsuite-ready -- the breakage demonstrated depends on the stack layout, and it only "works" at -O0. HTH, Julian
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 02c40fdc660..909a86795e0 100644 --- 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); + } + else if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))) { tree type = TREE_TYPE (decl); tree ptr = gfc_conv_descriptor_data_get (decl); diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 new file mode 100644 index 00000000000..2dd1a6fa5b6 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 @@ -0,0 +1,51 @@ +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) + + ! FIXME: This warning is emitted on the wrong line number. + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 36 } + !$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) + call acc_copyout(myvar) + call acc_copyout(tarr) + + 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