From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id 8108C3857C6E; Fri, 17 Jul 2020 11:16:23 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 8108C3857C6E Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Thomas_Schwinge@mentor.com IronPort-SDR: CqMAJm3TG0jXglYvG5FKBIa3AK73JvMWi6FhvGNvGTsohKYZbyys//q8kuRwEGstATeZlXEDLR PEX/vxVioXUhOdZfJCJtHwO4478lfon6RV7P/6p7yJfB7ZdES0Lx4Eb9SGv85JuViMWrKE5mJs q0KpPBUr8iHI4uiLoiFSeexKPIh/V2z+/vIf47zA12YTdVnyWTnEtWy3uq5K2bZF4bZczf4VlJ Oerv7oQLalUFMk0S65ziK1mWFWbroYsbawBP7LUyWbYOEZaJedDMobWkA0CMbfx0bsSckuUtu4 Qrs= X-IronPort-AV: E=Sophos;i="5.75,362,1589270400"; d="scan'208,223";a="53163639" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 17 Jul 2020 03:16:20 -0800 IronPort-SDR: NVh4Qxf5jQmVeOao7KVkArJnj8k5esD8fbRg1n+daaZV0DApD7z0rQa94HVuYeJSKDqf44YRXk 9/5fgr9acCFBh2evP3TodIHKVrirUz53MlcThbP+c7Ki5A2ooFXw1XEay+zY97jd7sExZXmogn fKK4fV8BfAAGkLNZnMl/BgwxQd31JArgWeB90qh/pWLEfqMHBp1fKV5Fls6Da52qWg+ZiPcTK2 AelmJOwCuzC0doSRJbsDwaXyk+S5LH8PwGbzXmiqIsqfbnZNlq3R/zlszb9fv0qrQ+jGZfuXFl kj0= From: Thomas Schwinge To: Julian Brown , Tobias Burnus CC: , Jakub Jelinek , , Subject: Re: [PATCH 8/9] [OpenACC] Fix standalone attach for Fortran assumed-shape array pointers In-Reply-To: <87zh81rs11.fsf@euler.schwinge.homeip.net> References: <0193ff08d4a4a2b6ca86c7a891b8ff35203fa440.1592343757.git.julian@codesourcery.com> <877dv6tj86.fsf@euler.schwinge.homeip.net> <87zh81rs11.fsf@euler.schwinge.homeip.net> User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.3 (x86_64-pc-linux-gnu) Date: Fri, 17 Jul 2020 13:16:11 +0200 Message-ID: <877dv2s878.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="=-=-=" X-Spam-Status: No, score=-10.3 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , X-List-Received-Date: Fri, 17 Jul 2020 11:16:26 -0000 --=-=-= Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable Hi Julian, Tobias! On 2020-07-15T12:28:42+0200, Thomas Schwinge wrot= e: > On 2020-07-14T13:43:37+0200, I wrote: >> On 2020-06-16T15:39:44-0700, Julian Brown wrot= e: >>> 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 >> , >> you (Julian) correctly supposed in >> , 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), an= d >> 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", > . > > 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_o= mp_clauses *clauses, >>> } >>> } >>> if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)) >>> - && n->u.map_op !=3D OMP_MAP_ATTACH >>> - && n->u.map_op !=3D OMP_MAP_DETACH) >>> + && (n->u.map_op =3D=3D OMP_MAP_ATTACH >>> + || n->u.map_op =3D=3D OMP_MAP_DETACH)) >>> + { >>> + tree type =3D TREE_TYPE (decl); >>> + tree data =3D gfc_conv_descriptor_data_get (decl); >>> + if (present) >>> + data =3D gfc_build_cond_assign_expr (block, present= , >>> + data, >>> + null_pointer_nod= e); >>> + tree ptr >>> + =3D fold_convert (build_pointer_type (char_type_nod= e), >>> + data); >>> + ptr =3D 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) =3D ptr; >>> + node2 =3D build_omp_clause (input_location, OMP_CLAUS= E_MAP); >>> + OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET); >>> + OMP_CLAUSE_DECL (node2) =3D decl; >>> + OMP_CLAUSE_SIZE (node2) =3D TYPE_SIZE_UNIT (type); >>> + node3 =3D build_omp_clause (input_location, OMP_CLAUS= E_MAP); >>> + if (n->u.map_op =3D=3D OMP_MAP_ATTACH) >>> + { >>> + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH); >>> + n->u.map_op =3D OMP_MAP_ALLOC; >>> + } >>> + else /* OMP_MAP_DETACH. */ >>> + { >>> + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH); >>> + n->u.map_op =3D OMP_MAP_RELEASE; >>> + } >>> + OMP_CLAUSE_DECL (node3) =3D data; >>> + OMP_CLAUSE_SIZE (node3) =3D 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 =3D=3D false', if that makes sense, or a= ny >> 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=3D1,10 >>> + myvar%arr1(i) =3D 0 >>> + myvar%arr2(i) =3D 0 >>> + tarr(i) =3D 0 >>> + end do >>> + >>> + call acc_copyin(myvar) >>> + call acc_copyin(myvar%arr2) >>> + call acc_copyin(tarr) >>> + >>> + myptr =3D> 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" "" { targe= t openacc_nvidia_accel_selected } 36 } >> >> ... don't forget to adjust "36" here. ;-) >> >>> + !$acc serial present(myvar%arr2) >>> + do i=3D1,10 >>> + myvar%arr1(i) =3D i >>> + myvar%arr2(i) =3D i >>> + end do >>> + myptr(3) =3D 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 =3D> tarr', and thus... >>> + >>> + call acc_copyout(myvar%arr2) >>> + call acc_copyout(myvar) >>> + call acc_copyout(tarr) >>> + >>> + do i=3D1,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) =3D=3D 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 =3D> 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 =3D> 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) =3D=3D 1' myptr =3D> tarr !$acc enter data attach(myptr) ! 'rc(tarr) =3D=3D 2'! (not intended by = the user) !$acc enter data attach(myptr) ! 'rc(tarr) =3D=3D 3'! (not intended by = the user) [...] call acc_copyout(tarr) ! won't copyout, because still 'rc(tarr) =3D 2'!= (not intended by the user) if (acc_is_present(tarr)) stop 12 ! fails Ugh. :-( Or am I confused now? Gr=C3=BC=C3=9Fe Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstra=C3=9Fe 201, 80634 M=C3=BCnch= en / Germany Registergericht M=C3=BCnchen HRB 106955, Gesch=C3=A4ftsf=C3=BChrer: Thomas = Heurung, Alexander Walter --=-=-= Content-Type: text/x-diff Content-Disposition: inline; filename="0001-WIP-Problems-with-openacc-Fix-standalone-attach-for-.patch" >From 4fa4979da3de6d15d5a39b77fdeb6b5aadec0f10 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Fri, 17 Jul 2020 13:09:58 +0200 Subject: [PATCH] WIP Problems with "openacc: Fix standalone attach for Fortran assumed-shape array pointers" --- .../gfortran.dg/goacc/attach-descriptor.f90 | 11 ++++++++++- .../libgomp.oacc-fortran/attach-descriptor-1.f90 | 3 +++ ...descriptor-1.f90 => attach-descriptor-1_.f90} | 8 ++++++++ ...escriptor-1.f90 => attach-descriptor-1__.f90} | 16 +++++++++++++++- 4 files changed, 36 insertions(+), 2 deletions(-) copy libgomp/testsuite/libgomp.oacc-fortran/{attach-descriptor-1.f90 => attach-descriptor-1_.f90} (79%) copy libgomp/testsuite/libgomp.oacc-fortran/{attach-descriptor-1.f90 => attach-descriptor-1__.f90} (62%) diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 index 9ca36f770c7..454ef9cccf3 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 @@ -12,7 +12,16 @@ program att !$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 omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(alloc:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) 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 omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(release:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) 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\\(release:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize;$" 1 "original" } } + !TODO See how in the 'gimple' dump, 'detach' is turned into 'force_detach', and 'release' into 'delete' -- but is the latter actually correct? (See 'libgomp.oacc-fortran/attach-descriptor-1__.f90'.) +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) 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..99c1e787de6 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 @@ -41,8 +41,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-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1_.f90 similarity index 79% copy from libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 copy to libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1_.f90 index 5d79cbc14fc..2e2e1267660 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1_.f90 @@ -26,6 +26,7 @@ program att 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. @@ -39,10 +40,17 @@ program att !$acc end serial !$acc exit data detach(myvar%arr2, myptr) + !!$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 ! fails + + !TODO Have to stop, have copied out device pointers. + stop do i=1,10 if (myvar%arr1(i) .ne. i) stop 1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1__.f90 similarity index 62% copy from libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 copy to libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1__.f90 index 5d79cbc14fc..6786f32852b 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1__.f90 @@ -27,6 +27,7 @@ program att 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 } 38 } @@ -38,11 +39,24 @@ program att myptr(3) = 99 !$acc end serial - !$acc exit data detach(myvar%arr2, myptr) + !$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 ! fails 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 -- 2.17.1 --=-=-=--