From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id AF401386F037; Mon, 27 Jul 2020 14:34:20 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org AF401386F037 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: H1IdTyLtHgvBQVxuKWjbijZMzOYIvOpW37gik+E/kZ96cVuA74WEWhJYUUDINf0vXjF1mUAKim QSnqjzYjx3dr4GD12HYKOG2toDjgSQX0dpNitEe0sfE+ar0WhqDa9VeYztfs4Yrr/bOMuHg3E9 RAvv6XJUIU21na/5yVjSrTs64J2Ys0MqKsNelAZo16oftNi8ugyKIwjGiN1tVKMR+LGizKevKb FYqMxEcludg6ywqi9liCJ/q36jWdOtiMFsZ3LZ704B7KLmTQRIT8wSCpvZjJEc/tZmOZrdIsRt JRk= X-IronPort-AV: E=Sophos;i="5.75,402,1589270400"; d="scan'208,223";a="51325496" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 27 Jul 2020 06:34:18 -0800 IronPort-SDR: siCYOaR4MvKRgDoH6PVnpnG857KqJQ2/yIopjzpNER2ExJo1ncXs5AGOWyko7WrBahsF5LWXa9 ao3dt65TBVsacI6+RHS5AJtmRh/uivH3x6gFnn22gIl9flxF5WzSUVuzJQQk3UyDZgOXq9FiJ4 iq7LdrgVFZ0YL2fVgzrteMsW0WitJn/ttODzmFcFEK59eHm/LhltFOcf2bHUnLBQUGYj9wYogR XK9GHvI1IziWWOkZUZbJSdTQfdJLLLUzc6lox4Tlqk/Io0xo25KgAb8kfVi2HJYbxPLT2fNIX+ Kjw= Date: Mon, 27 Jul 2020 15:33:41 +0100 From: Julian Brown To: Thomas Schwinge CC: Tobias Burnus , , Jakub Jelinek , , Subject: Re: [PATCH 8/9] [OpenACC] Fix standalone attach for Fortran assumed-shape array pointers Message-ID: <20200727153341.00a59652@squid.athome> In-Reply-To: <877dv2s878.fsf@euler.schwinge.homeip.net> References: <0193ff08d4a4a2b6ca86c7a891b8ff35203fa440.1592343757.git.julian@codesourcery.com> <877dv6tj86.fsf@euler.schwinge.homeip.net> <87zh81rs11.fsf@euler.schwinge.homeip.net> <877dv2s878.fsf@euler.schwinge.homeip.net> Organization: Mentor Graphics X-Mailer: Claws Mail 3.17.5 (GTK+ 2.24.32; x86_64-pc-linux-gnu) MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="MP_/vlHxBMyYT332=.BWLvEZnCK" X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-11.9 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: Mon, 27 Jul 2020 14:34:24 -0000 --MP_/vlHxBMyYT332=.BWLvEZnCK Content-Type: text/plain; charset="US-ASCII" Content-Transfer-Encoding: 7bit Content-Disposition: inline On Fri, 17 Jul 2020 13:16:11 +0200 Thomas Schwinge wrote: > Hi Julian, Tobias! > > On 2020-07-15T12:28:42+0200, Thomas Schwinge > wrote: > > On 2020-07-14T13:43:37+0200, I wrote: > >> On 2020-06-16T15:39:44-0700, Julian Brown > >> 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 > >> , > >> 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), 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", > > . > > > > 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 --MP_/vlHxBMyYT332=.BWLvEZnCK Content-Type: text/x-patch Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="0001-openacc-No-attach-detach-present-release-mappings-fo.patch" >From d53e4f1cd450062163e7e96a469c2f56cfac65ee Mon Sep 17 00:00:00 2001 From: Julian Brown Date: Mon, 27 Jul 2020 06:29:02 -0700 Subject: [PATCH] openacc: No attach/detach present/release mappings for array descriptors Standalone attach and detach clauses should not create present/release mappings for Fortran array descriptors (e.g. used when we have a pointer to an array), both because it is unnecessary and because those mappings will be incorrectly subject to reference counting. Simply omitting the mappings means we just use GOMP_MAP_TO_PSET and GOMP_MAP_{ATTACH,DETACH} mappings for array descriptors. That requires a tweak in gimplify.c, since we may now see GOMP_MAP_TO_PSET without a preceding data-movement mapping. The new attach-descriptor-4.f90 test relies on the checking performed by the patch "Refuse update/copyout for blocks with attached pointers". 2020-07-27 Julian Brown Thomas Schwinge gcc/fortran/ * trans-openmp.c (gfc_trans_omp_clauses): Don't create present/release mappings for array descriptors. gcc/ * gimplify.c (gimplify_omp_target_update): Allow GOMP_MAP_TO_PSET without a preceding data-movement mapping. gcc/testsuite/ * gfortran.dg/goacc/attach-descriptor.f90: Update pattern output. Add scanning of gimplify dump. libgomp/ * testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90: Don't run for shared-memory devices. Add more checking. * testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90: New test. * testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90: New test. Co-Authored-By: Thomas Schwinge --- gcc/fortran/trans-openmp.c | 44 +++++++----- gcc/gimplify.c | 3 +- .../gfortran.dg/goacc/attach-descriptor.f90 | 17 ++++- .../attach-descriptor-1.f90 | 6 +- .../attach-descriptor-3.f90 | 68 +++++++++++++++++++ .../attach-descriptor-4.f90 | 61 +++++++++++++++++ 6 files changed, 177 insertions(+), 22 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90 diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index d12d7fbddac..1a8f3277de3 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2718,23 +2718,6 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type); node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP); - if (n->u.map_op == OMP_MAP_ATTACH) - { - /* Standalone attach clauses used with arrays with - descriptors must copy the descriptor to the target, - else they won't have anything to perform the - attachment onto (see OpenACC 2.6, "2.6.3. Data - Structures with Pointers"). */ - OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC); - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH); - } - else if (n->u.map_op == OMP_MAP_DETACH) - { - OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE); - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH); - } - else - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER); if (present) { ptr = gfc_conv_descriptor_data_get (decl); @@ -2748,6 +2731,33 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (decl); OMP_CLAUSE_SIZE (node3) = size_int (0); + if (n->u.map_op == OMP_MAP_ATTACH) + { + /* Standalone attach clauses used with arrays with + descriptors must copy the descriptor to the target, + else they won't have anything to perform the + attachment onto (see OpenACC 2.6, "2.6.3. Data + Structures with Pointers"). */ + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH); + /* We don't want to map PTR at all in this case, so + delete its node and shuffle the others down. */ + node = node2; + node2 = node3; + node3 = NULL; + goto finalize_map_clause; + } + else if (n->u.map_op == OMP_MAP_DETACH) + { + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH); + /* Similarly to above, we don't want to unmap PTR + here. */ + node = node2; + node2 = node3; + node3 = NULL; + goto finalize_map_clause; + } + else + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER); /* We have to check for n->sym->attr.dimension because of scalar coarrays. */ diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 15dfee903ab..f4c31d2870d 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -13013,8 +13013,9 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p) OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE); have_clause = true; break; - case GOMP_MAP_POINTER: case GOMP_MAP_TO_PSET: + break; + case GOMP_MAP_POINTER: /* TODO PR92929: we may see these here, but they'll always follow one of the clauses above, and will be handled by libgomp as one group, so no handling required here. */ diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 index 9ca36f770c7..373bdcb2114 100644 --- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 @@ -1,4 +1,4 @@ -! { dg-additional-options "-fdump-tree-original" } +! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } program att implicit none @@ -11,8 +11,19 @@ program att integer, pointer :: myptr(:) !$acc enter data attach(myvar%arr2, myptr) -! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(alloc:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } } !$acc exit data detach(myvar%arr2, myptr) -! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(release:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } } + + ! Test valid usage and processing of the finalize clause. + !$acc exit data detach(myvar%arr2, myptr) finalize +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize;$" 1 "original" } } + ! For array-descriptor detaches, we no longer generate a "release" mapping + ! for the pointed-to data for gimplify.c to turn into "delete". Make sure + ! the mapping still isn't there. +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } } + end program att diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 index 5d79cbc14fc..9f159fa3b75 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 @@ -1,4 +1,5 @@ ! { dg-do run } +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } program att use openacc @@ -29,7 +30,7 @@ program att !$acc enter data attach(myvar%arr2, myptr) ! FIXME: This warning is emitted on the wrong line number. - ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 38 } + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 39 } !$acc serial present(myvar%arr2) do i=1,10 myvar%arr1(i) = i @@ -41,8 +42,11 @@ program att !$acc exit data detach(myvar%arr2, myptr) call acc_copyout(myvar%arr2) + if (acc_is_present(myvar%arr2)) stop 10 call acc_copyout(myvar) + if (acc_is_present(myvar)) stop 11 call acc_copyout(tarr) + if (acc_is_present(tarr)) stop 12 do i=1,10 if (myvar%arr1(i) .ne. i) stop 1 diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90 new file mode 100644 index 00000000000..f0e57b47453 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-3.f90 @@ -0,0 +1,68 @@ +! { dg-do run } +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } + +program att + use openacc + implicit none + type t + integer :: arr1(10) + integer, allocatable :: arr2(:) + end type t + integer :: i + type(t) :: myvar + integer, target :: tarr(10) + integer, pointer :: myptr(:) + + allocate(myvar%arr2(10)) + + do i=1,10 + myvar%arr1(i) = 0 + myvar%arr2(i) = 0 + tarr(i) = 0 + end do + + call acc_copyin(myvar) + call acc_copyin(myvar%arr2) + call acc_copyin(tarr) + + myptr => tarr + + !$acc enter data attach(myvar%arr2, myptr) + !$acc enter data attach(myvar%arr2, myptr) + + ! FIXME: This warning is emitted on the wrong line number. + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 40 } + !$acc serial present(myvar%arr2) + do i=1,10 + myvar%arr1(i) = i + myvar%arr2(i) = i + end do + myptr(3) = 99 + !$acc end serial + + !$acc exit data detach(myvar%arr2, myptr) finalize + + if (.not. acc_is_present(myvar%arr2)) stop 10 + if (.not. acc_is_present(myvar)) stop 11 + if (.not. acc_is_present(tarr)) stop 12 + + call acc_copyout(myvar%arr2) + if (acc_is_present(myvar%arr2)) stop 20 + if (.not. acc_is_present(myvar)) stop 21 + if (.not. acc_is_present(tarr)) stop 22 + call acc_copyout(myvar) + if (acc_is_present(myvar%arr2)) stop 30 + if (acc_is_present(myvar)) stop 31 + if (.not. acc_is_present(tarr)) stop 32 + call acc_copyout(tarr) + if (acc_is_present(myvar%arr2)) stop 40 + if (acc_is_present(myvar)) stop 41 + if (acc_is_present(tarr)) stop 42 + + do i=1,10 + if (myvar%arr1(i) .ne. i) stop 1 + if (myvar%arr2(i) .ne. i) stop 2 + end do + if (tarr(3) .ne. 99) stop 3 + +end program att diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90 new file mode 100644 index 00000000000..9dbf53d0213 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-4.f90 @@ -0,0 +1,61 @@ +! { dg-do run } +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } + +program att + use openacc + implicit none + type t + integer :: arr1(10) + integer, allocatable :: arr2(:) + end type t + integer :: i + type(t) :: myvar + integer, target :: tarr(10) + integer, pointer :: myptr(:) + + allocate(myvar%arr2(10)) + + do i=1,10 + myvar%arr1(i) = 0 + myvar%arr2(i) = 0 + tarr(i) = 0 + end do + + call acc_copyin(myvar) + call acc_copyin(myvar%arr2) + call acc_copyin(tarr) + + myptr => tarr + + !$acc enter data attach(myvar%arr2, myptr) + !$acc enter data attach(myvar%arr2, myptr) + + ! FIXME: This warning is emitted on the wrong line number. + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 40 } + !$acc serial present(myvar%arr2) + do i=1,10 + myvar%arr1(i) = i + myvar%arr2(i) = i + end do + myptr(3) = 99 + !$acc end serial + + !$acc exit data detach(myvar%arr2, myptr) + + call acc_copyout(myvar%arr2) + ! { dg-output ".*copyout of block \\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\] with attached pointers(\n|\r\n|\r)+" } + if (acc_is_present(myvar%arr2)) stop 10 + call acc_copyout(myvar) + if (acc_is_present(myvar)) stop 11 + call acc_copyout(tarr) + if (acc_is_present(tarr)) stop 12 + + do i=1,10 + if (myvar%arr1(i) .ne. i) stop 1 + if (myvar%arr2(i) .ne. i) stop 2 + end do + if (tarr(3) .ne. 99) stop 3 + +end program att + +! { dg-shouldfail "" } -- 2.23.0 --MP_/vlHxBMyYT332=.BWLvEZnCK--