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 39AF03857C4A for ; Thu, 16 Jul 2020 08:36:31 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 39AF03857C4A 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: fK5KTL6kfE2t0NrSvkAzydROjct5N9iObfz2Uj1zNq+V8b8iZsA9iMY4oNqCqR/8s050wFF/Yl /O4T+mU7efUMdX8maj5XojfyMdN2idPV6bfUxDNhCbY3YzjBazcS/XdaV4+RkZaPO7Z+EL46LI YlK9MWAunf9Fkt4I7tivxO48YKU2/7jkOeULpoyTlvJcGl1clbxjstD1pr8GsGlI/102ZjXH+u ops4hpXF8d9tM9LlLXec2XC3zkuoqMmoYf/xzTxwYFn6/ujY955PTFuomoXWysZtEGjrWPJ61+ 5Dg= X-IronPort-AV: E=Sophos;i="5.75,358,1589270400"; d="scan'208,223";a="50965065" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 16 Jul 2020 00:36:10 -0800 IronPort-SDR: zssxtezDhQfeuDxxiyvPv6Auddiy8eQuFCpcXp3CjfSKdzMROkufUGq9PBI82eC1rkPntFh9Pr saSzTIfv9vOKVZjv9PxoEXdgkqBeP7L3UmwQATJk7d8BNJTmfwN9cwmJiQpkzTWo8L7Wz+HDMl OJQNhDkWk7GIWzlnTCPTLLqykAXr1R0kuwqynKDywTV8tW4lfh8BXxCEe05euuiLYhqZWT6TEw d39nrDIpORjFmZyxWfQ2rtp5hJ0gZATJoBAWKpP46+4JEcttrLypr3IePHjE3mqqWap+e+Yupg b4w= From: Thomas Schwinge To: Julian Brown CC: Jakub Jelinek , Subject: Re: OpenACC 'attach'/'detach' has no business affecting user-visible reference counting In-Reply-To: <20200618192157.553dea83@squid.athome> References: <65540b92dff74db1f15af930f87f7096d03e7efe.1576648001.git.julian@codesourcery.com> <20200605213108.7f2e4807@squid.athome> <87k10gr06m.fsf@euler.schwinge.homeip.net> <20200618192157.553dea83@squid.athome> User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.3 (x86_64-pc-linux-gnu) Date: Thu, 16 Jul 2020 10:35:53 +0200 Message-ID: <87sgdrsvpy.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="=-=-=" X-Spam-Status: No, score=-10.4 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: Thu, 16 Jul 2020 08:36:34 -0000 --=-=-= Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable Hi Julian! On 2020-06-18T19:21:57+0100, Julian Brown wrote: > On Tue, 9 Jun 2020 12:41:21 +0200 > Thomas Schwinge wrote: >> On 2020-06-05T21:31:08+0100, Julian Brown >> wrote: >> > On Fri, 5 Jun 2020 13:17:09 +0200 >> > Thomas Schwinge wrote: >> >> On 2019-12-17T21:03:47-0800, Julian Brown >> >> wrote: >> >> > This part contains the libgomp runtime support for the >> >> > GOMP_MAP_ATTACH and GOMP_MAP_DETACH mapping kinds >> >> >> >> > --- a/libgomp/target.c >> >> > +++ b/libgomp/target.c >> >> >> >> > @@ -1203,6 +1211,32 @@ gomp_map_vars_internal (struct gomp_device_d= escr *devicep, >> >> >> >> > + case GOMP_MAP_ATTACH: >> >> > + { >> >> > + cur_node.host_start =3D (uintptr_t) hostaddrs[i]; >> >> > + cur_node.host_end =3D cur_node.host_start + sizeo= f (void *); >> >> > + splay_tree_key n =3D splay_tree_lookup (mem_map, = &cur_node); >> >> > + if (n !=3D NULL) >> >> > + { >> >> > + tgt->list[i].key =3D n; >> >> > + tgt->list[i].offset =3D cur_node.host_start -= n->host_start; >> >> > + tgt->list[i].length =3D n->host_end - n->host= _start; >> >> > + tgt->list[i].copy_from =3D false; >> >> > + tgt->list[i].always_copy_from =3D false; >> >> > + tgt->list[i].do_detach >> >> > + =3D (pragma_kind !=3D GOMP_MAP_VARS_OPENACC= _ENTER_DATA); >> >> > + n->refcount++; >> >> > + } >> >> > + else >> >> > + { >> >> > + gomp_mutex_unlock (&devicep->lock); >> >> > + gomp_fatal ("outer struct not mapped for atta= ch"); >> >> > + } >> >> > + gomp_attach_pointer (devicep, aq, mem_map, n, >> >> > + (uintptr_t) hostaddrs[i], si= zes[i], >> >> > + cbufp); >> >> > + continue; >> >> > + } >> >> >> >> For the OpenACC runtime API 'acc_attach' etc. routines they don't, >> >> so what's the conceptual reason that for the corresponding OpenACC >> >> directive variants, 'GOMP_MAP_ATTACH' etc. here participate in >> >> reference counting ('n->refcount++' above)? I understand OpenACC >> >> 'attach'/'detach' clauses to be simple "executable clauses", which >> >> just update some values somewhere (say, like >> >> 'GOMP_MAP_ALWAYS_POINTER'), but they don't alter any mapping state, >> >> thus wouldn't appear to need reference counting? >> > >> > IIUC, n->refcount is not directly the "structural reference count" >> > as seen at source level, but rather counts the number of >> > target_var_descs in the lists appended to each target_mem_desc -- >> > and GOMP_MAP_ATTACH have variable entries in those lists. >> >> That may be OK if that's purely an implementation detail that isn't >> visible to the user, however: >> >> > That's not the case for the API >> > routines. >> >> As I had mentioned, the problem is: in contrast to 'acc_attach', an >> OpenACC 'enter data' directive with 'attach' clause currently uses >> this same reference-counted code path, and thus such an 'attach' >> without corresponding 'detach' inhibits unmapping; [...] > > The attached patch stops attach/detach operations from affecting > reference counts (either structured or dynamic). This isn't as invasive > as I'd imagined: we can extend the use of the "do_detach" flag in > target_mem_descs' variable lists to mark mappings that correspond to > attach operations, then use that flag to avoid refcount > increment/decrements. Thanks, ACK. > (The flag should possibly be renamed now.) How about: - /* True if variable should be detached at end of region. */ - bool do_detach; + /* True if this is for OpenACC 'attach'. */ + bool is_attach; (Changing that similarly is obvious/pre-approved.) > Tested with offloading to NVPTX. OK? I've adjusted the patch for current GCC sources, and did some further changes/cleanup; see below, and attached "[OpenACC] Deep copy attach/detach should not affect reference counts". If you're happy with that, that's OK for master and releases/gcc-10 (once un-frozen) branches. > --- a/libgomp/oacc-mem.c > +++ b/libgomp/oacc-mem.c > @@ -1131,7 +1134,9 @@ goacc_enter_data_internal (struct gomp_device_descr= *acc_dev, size_t mapnum, > if (tgt->list[j].key =3D=3D n) > { > for (size_t k =3D 0; k < groupnum; k++) > - if (j + k < tgt->list_count && tgt->list[j + k].key) > + if (j + k < tgt->list_count > + && tgt->list[j + k].key > + && !tgt->list[j + k].do_detach) > { > tgt->list[j + k].key->refcount++; > tgt->list[j + k].key->dynamic_refcount++; > @@ -1156,7 +1161,7 @@ goacc_enter_data_internal (struct gomp_device_descr= *acc_dev, size_t mapnum, > for (size_t j =3D 0; j < tgt->list_count; j++) > { > n =3D tgt->list[j].key; > - if (n) > + if (n && !tgt->list[j].do_detach) > n->dynamic_refcount++; > } > } If I understand correctly, relatedly, we can also "strengthen" the 'is_tgt_unmapped' checking (nowadays centralized in 'goacc_exit_datum_1') by excluding any 'do_detach' ones from '++num_mappings'. Done. > --- a/libgomp/target.c > +++ b/libgomp/target.c > @@ -382,7 +382,7 @@ gomp_map_vars_existing (struct gomp_device_descr *dev= icep, > (void *) newn->host_start, > newn->host_end - newn->host_start, cbuf); > > - if (oldn->refcount !=3D REFCOUNT_INFINITY) > + if (oldn->refcount !=3D REFCOUNT_INFINITY && kind !=3D GOMP_MAP_ATTACH= ) > oldn->refcount++; > } That's always-true. Removed. > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/attach-detach-rc-1.c > @@ -0,0 +1,50 @@ > +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=3D1" } } */ > + > +#include > +#include > + > +#define N 1024 > + > +struct mystr { > + int pad; > + int *data; > +}; The 'pad' is no longer needed with PR95270 "OpenACC 'enter data attach' looks up target memory object displaced by pointer size" fixed. > +[...] > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/attach-detach-rc-2.c > @@ -0,0 +1,4 @@ > +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=3D1" } } */ > +/* { dg-additional-options "-DATTACH_VIA_DIRECTIVE" } */ > + > +#include "attach-detach-rc-1.c" I've merged/extended 'libgomp.oacc-c-c++-common/attach-detach-rc-1.c', 'libgomp.oacc-c-c++-common/attach-detach-rc-2.c' into 'libgomp.oacc-c-c++-common/mdc-refcount-1.c', and further added 'libgomp.oacc-c-c++-common/mdc-refcount-2.c', and 'libgomp.oacc-c-c++-common/mdc-refcount-3.c'. 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-OpenACC-Deep-copy-attach-detach-should-not-affect-re.patch" >From 3b1262da8922df1321ab982744ac48334b2279da Mon Sep 17 00:00:00 2001 From: Julian Brown Date: Thu, 18 Jun 2020 05:11:08 -0700 Subject: [PATCH] [OpenACC] Deep copy attach/detach should not affect reference counts TODO Some rationale. TODO Update libgomp/ * oacc-mem.c (goacc_enter_data_internal): Don't affect reference counts for attach mappings. (goacc_exit_data_internal): Don't affect reference counts for detach mappings. * target.c (gomp_map_vars_existing): Don't affect reference counts for attach mappings. (gomp_map_vars_internal): Set do_detach flag unconditionally to mark attach mappings. (gomp_unmap_vars_internal): Use above flag to prevent affecting reference count for attach mappings. * testsuite/libgomp.oacc-c-c++-common/attach-detach-rc-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/attach-detach-rc-2.c: Likewise. * testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90: Mark test as shouldfail. * testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Adjust to fail gracefully in no-finalize mode. Co-authored-by: Thomas Schwinge --- libgomp/oacc-mem.c | 40 +++--- libgomp/target.c | 12 +- .../mdc-refcount-1.c | 60 +++++++++ .../mdc-refcount-2.c | 123 ++++++++++++++++++ .../mdc-refcount-3.c | 86 ++++++++++++ .../deep-copy-6-no_finalize.F90 | 9 +- .../libgomp.oacc-fortran/deep-copy-6.f90 | 8 +- 7 files changed, 318 insertions(+), 20 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/mdc-refcount-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/mdc-refcount-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/mdc-refcount-3.c diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 41548f75e72c..0fa6597aaf1b 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -669,6 +669,9 @@ static void goacc_exit_datum_1 (struct gomp_device_descr *acc_dev, void *h, size_t s, unsigned short kind, splay_tree_key n, goacc_aq aq) { + assert (kind != GOMP_MAP_DETACH + && kind != GOMP_MAP_FORCE_DETACH); + if ((uintptr_t) h < n->host_start || (uintptr_t) h + s > n->host_end) { size_t host_size = n->host_end - n->host_start; @@ -678,8 +681,7 @@ goacc_exit_datum_1 (struct gomp_device_descr *acc_dev, void *h, size_t s, } bool finalize = (kind == GOMP_MAP_FORCE_FROM - || kind == GOMP_MAP_DELETE - || kind == GOMP_MAP_FORCE_DETACH); + || kind == GOMP_MAP_DELETE); assert (n->refcount != REFCOUNT_LINK); if (n->refcount != REFCOUNT_INFINITY @@ -727,7 +729,8 @@ goacc_exit_datum_1 (struct gomp_device_descr *acc_dev, void *h, size_t s, zero. Otherwise (e.g. for a 'GOMP_MAP_STRUCT' mapping with multiple members), fall back to skipping the test. */ for (size_t l_i = 0; l_i < n->tgt->list_count; ++l_i) - if (n->tgt->list[l_i].key) + if (n->tgt->list[l_i].key + && !n->tgt->list[l_i].do_detach) ++num_mappings; bool is_tgt_unmapped = gomp_remove_var (acc_dev, n); assert (is_tgt_unmapped || num_mappings > 1); @@ -1137,12 +1140,15 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, void *h = hostaddrs[i]; size_t s = sizes[i]; - /* A standalone attach clause. */ if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH) - gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, - (uintptr_t) h, s, NULL); - - goacc_map_var_existing (acc_dev, h, s, n); + { + gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, + (uintptr_t) h, s, NULL); + /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic + reference counts ('n->refcount', 'n->dynamic_refcount'). */ + } + else + goacc_map_var_existing (acc_dev, h, s, n); } else if (n && groupnum > 1) { @@ -1170,7 +1176,9 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, list, and increment the refcounts for each item in that group. */ for (size_t k = 0; k < groupnum; k++) - if (j + k < tgt->list_count && tgt->list[j + k].key) + if (j + k < tgt->list_count + && tgt->list[j + k].key + && !tgt->list[j + k].do_detach) { tgt->list[j + k].key->refcount++; tgt->list[j + k].key->dynamic_refcount++; @@ -1204,7 +1212,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, for (size_t j = 0; j < tgt->list_count; j++) { n = tgt->list[j].key; - if (n) + if (n && !tgt->list[j].do_detach) n->dynamic_refcount++; } } @@ -1270,14 +1278,10 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, case GOMP_MAP_POINTER: case GOMP_MAP_DELETE: case GOMP_MAP_RELEASE: - case GOMP_MAP_DETACH: - case GOMP_MAP_FORCE_DETACH: { struct splay_tree_key_s cur_node; size_t size; - if (kind == GOMP_MAP_POINTER - || kind == GOMP_MAP_DETACH - || kind == GOMP_MAP_FORCE_DETACH) + if (kind == GOMP_MAP_POINTER) size = sizeof (void *); else size = sizes[i]; @@ -1300,6 +1304,12 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, 'GOMP_MAP_STRUCT's anymore. */ break; + case GOMP_MAP_DETACH: + case GOMP_MAP_FORCE_DETACH: + /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic + reference counts ('n->refcount', 'n->dynamic_refcount'). */ + break; + default: gomp_fatal (">>>> goacc_exit_data_internal UNHANDLED kind 0x%.2x", kind); diff --git a/libgomp/target.c b/libgomp/target.c index 478909e3b275..0358864608a2 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1095,9 +1095,10 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, tgt->list[i].length = n->host_end - n->host_start; tgt->list[i].copy_from = false; tgt->list[i].always_copy_from = false; - tgt->list[i].do_detach - = (pragma_kind != GOMP_MAP_VARS_ENTER_DATA); - n->refcount++; + tgt->list[i].do_detach = true; + /* OpenACC 'attach'/'detach' doesn't affect + structured/dynamic reference counts ('n->refcount', + 'n->dynamic_refcount'). */ } else { @@ -1448,6 +1449,11 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom, if (k == NULL) continue; + /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic reference + counts ('n->refcount', 'n->dynamic_refcount'). */ + if (tgt->list[i].do_detach) + continue; + bool do_unmap = false; if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY) k->refcount--; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/mdc-refcount-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/mdc-refcount-1.c new file mode 100644 index 000000000000..6170447e7d31 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/mdc-refcount-1.c @@ -0,0 +1,60 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include + +#define N 1024 + +struct mystr { + int *data; +}; + +static void +test (unsigned variant) +{ + int arr[N]; + struct mystr s; + + s.data = arr; + + acc_copyin (&s, sizeof (s)); + acc_create (s.data, N * sizeof (int)); + + for (int i = 0; i < 20; i++) + { + if ((variant + i) % 1) + { +#pragma acc enter data attach(s.data) + } + else + acc_attach ((void **) &s.data); + + if ((variant + i) % 2) + { +#pragma acc exit data detach(s.data) + } + else + acc_detach ((void **) &s.data); + } + + assert (acc_is_present (arr, N * sizeof (int))); + assert (acc_is_present (&s, sizeof (s))); + + acc_delete (arr, N * sizeof (int)); + + assert (!acc_is_present (arr, N * sizeof (int))); + + acc_copyout (&s, sizeof (s)); + + assert (!acc_is_present (&s, sizeof (s))); + assert (s.data == arr); +} + +int +main (int argc, char *argv[]) +{ + for (unsigned variant = 0; variant < 4; ++variant) + test (variant); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/mdc-refcount-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/mdc-refcount-2.c new file mode 100644 index 000000000000..2431a76a805c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/mdc-refcount-2.c @@ -0,0 +1,123 @@ +/* Verify that OpenACC 'attach'/'detach' doesn't interfere with reference + counting. */ + +#include +#include +#include + +/* Need to shared this (and, in particular, implicit '&data_work' in + 'attach'/'detach' clauses) between 'test' and 'test_'. */ +static unsigned char *data_work; + +static void test_(unsigned variant, + unsigned char *data, + void *data_d) +{ + assert(acc_is_present(&data_work, sizeof data_work)); + assert(data_work == data); + + acc_update_self(&data_work, sizeof data_work); + assert(data_work == data); + + if (variant & 1) + { +#pragma acc enter data attach(data_work) + } + else + acc_attach((void **) &data_work); + acc_update_self(&data_work, sizeof data_work); + assert(data_work == data_d); + + if (variant & 4) + { + if (variant & 2) + { // attach some more + data_work = data; + acc_attach((void **) &data_work); +#pragma acc enter data attach(data_work) + acc_attach((void **) &data_work); +#pragma acc enter data attach(data_work) +#pragma acc enter data attach(data_work) +#pragma acc enter data attach(data_work) + acc_attach((void **) &data_work); + acc_attach((void **) &data_work); +#pragma acc enter data attach(data_work) + } + else + {} + } + else + { // detach + data_work = data; + if (variant & 2) + { +#pragma acc exit data detach(data_work) + } + else + acc_detach((void **) &data_work); + acc_update_self(&data_work, sizeof data_work); + assert(data_work == data); + + // now not attached anymore + +#if 0 + if (TODO) + { + acc_detach(&data_work); //TODO PR95203 "libgomp: attach count underflow" + acc_update_self(&data_work, sizeof data_work); + assert(data_work == data); + } +#endif + } + + assert(acc_is_present(&data_work, sizeof data_work)); +} + +static void test(unsigned variant) +{ + const int size = sizeof (void *); + unsigned char *data = (unsigned char *) malloc(size); + assert(data); + void *data_d = acc_create(data, size); + assert(data_d); + assert(acc_is_present(data, size)); + + data_work = data; + + if (variant & 8) + { +#pragma acc data copyin(data_work) + test_(variant, data, data_d); + } + else + { + acc_copyin(&data_work, sizeof data_work); + test_(variant, data, data_d); + acc_delete(&data_work, sizeof data_work); + } +#if ACC_MEM_SHARED + assert(acc_is_present(&data_work, sizeof data_work)); +#else + assert(!acc_is_present(&data_work, sizeof data_work)); +#endif + data_work = NULL; + + assert(acc_is_present(data, size)); + acc_delete(data, size); + data_d = NULL; +#if ACC_MEM_SHARED + assert(acc_is_present(data, size)); +#else + assert(!acc_is_present(data, size)); +#endif + free(data); + data = NULL; +} + +int main() +{ + for (size_t i = 0; i < 16; ++i) + test(i); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/mdc-refcount-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/mdc-refcount-3.c new file mode 100644 index 000000000000..0f5e7becada8 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/mdc-refcount-3.c @@ -0,0 +1,86 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +/* Variant of 'deep-copy-7.c'. */ + +#include +#include +#include + +struct dc +{ + int a; + int *b; +}; + +int +main () +{ + int n = 100, i, j, k; + struct dc v = { .a = 3 }; + + v.b = (int *) malloc (sizeof (int) * n); + + for (k = 0; k < 16; k++) + { + /* Here, we do not explicitly copy the enclosing structure, but work + with fields directly. Make sure attachment counters and reference + counters work properly in that case. */ +#pragma acc enter data copyin(v.a, v.b[0:n]) // 1 + assert (acc_is_present (&v.b, sizeof v.b)); + assert (acc_is_present (v.b, sizeof (int) * n)); +#pragma acc enter data pcopyin(v.b[0:n]) // 2 +#pragma acc enter data pcopyin(v.b[0:n]) // 3 + +#pragma acc parallel loop present(v.a, v.b) + for (i = 0; i < n; i++) + v.b[i] = k + v.a + i; + + switch (k % 5) + { // All optional. + case 0: + break; + case 1: + ; //TODO PR95901 +#pragma acc exit data detach(v.b) finalize + break; + case 2: + ; //TODO PR95901 +#pragma acc exit data detach(v.b) + break; + case 3: + acc_detach_finalize ((void **) &v.b); + break; + case 4: + acc_detach ((void **) &v.b); + break; + } + assert (acc_is_present (&v.b, sizeof v.b)); + assert (acc_is_present (v.b, sizeof (int) * n)); + { // 3 + acc_delete (&v.b, sizeof v.b); + assert (acc_is_present (&v.b, sizeof v.b)); + acc_copyout (v.b, sizeof (int) * n); + assert (acc_is_present (v.b, sizeof (int) * n)); + } + { // 2 + acc_delete (&v.b, sizeof v.b); + assert (acc_is_present (&v.b, sizeof v.b)); + acc_copyout (v.b, sizeof (int) * n); + assert (acc_is_present (v.b, sizeof (int) * n)); + } + { // 1 + acc_delete (&v.b, sizeof v.b); + assert (!acc_is_present (&v.b, sizeof v.b)); + acc_copyout (v.b, sizeof (int) * n); + assert (!acc_is_present (v.b, sizeof (int) * n)); + } +#pragma acc exit data delete(v.a) + + for (i = 0; i < n; i++) + assert (v.b[i] == k + v.a + i); + + assert (!acc_is_present (&v, sizeof (v))); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90 index 038f04a3c37e..1daff2dadf11 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90 @@ -1,5 +1,12 @@ ! { dg-do run } -/* Nullify the 'finalize' clause. */ +/* Nullify the 'finalize' clause. + + That means, we do not detach properly, the host sees a device pointer, and + we fail as follows. + { dg-output "STOP 30(\n|\r\n|\r)+" { target { ! openacc_host_selected } } } + { dg-shouldfail "" { ! openacc_host_selected } } +*/ #define finalize #include "deep-copy-6.f90" + diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 index 6aab6a4a7633..94ddca3bce8e 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 @@ -12,11 +12,14 @@ program dtype end type mytype integer i - type(mytype) :: var + type(mytype), target :: var + integer, pointer :: hostptr(:) allocate(var%a(1:n)) allocate(var%b(1:n)) + hostptr => var%a + !$acc data copy(var) do i = 1, n @@ -49,6 +52,9 @@ program dtype !$acc end data + ! See 'deep-copy-6-no_finalize.F90'. + if (.not. associated(hostptr, var%a)) stop 30 + do i = 1,4 if (var%a(i) .ne. 0) stop 1 if (var%b(i) .ne. 0) stop 2 -- 2.27.0 --=-=-=--