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 C5CC538930DE for ; Fri, 24 Jul 2020 13:37:22 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org C5CC538930DE 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: wehvxtk5Ks1RcWWVlKBpE3dS+mS8CJS59yJbu+bx6dizapnvbG4AxrKBkPZud+gWZxr/0Pjk5P f/OO3BbFg+5NqDbRul6UGJOZLLxPHCzZHVPWkFBJAwC+auH4YXthPrOtZ/4PCJXAiJUEojdBhh OE1VqTri9NJVf1UNft48OZw9bT3bZ0DBOZw9PyQqrwyblthhH2FQPU8ZGBNm+kIVq4KPWtR6JD 8cIYsHTQtRshYPO6aQcKzm/ObnjsqCEkwVthnd3gA4vbKLaV3XPOPlogxjKVdEfxYDzp/7RmDK Dxo= X-IronPort-AV: E=Sophos;i="5.75,390,1589270400"; d="scan'208";a="53424354" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 24 Jul 2020 05:37:21 -0800 IronPort-SDR: KVvBYCnhzBrF3Lie+jmetoepiH2BSgtadTd2YzNVpTZFbdUdijGG51DmpW+FiLpgvx40drl2VH APtvf9FORenLO5j9GYZBcAkwvO4dLzE4RC3AU5XEOPq1oJLNbSjhvbP6kRQWpS+0GWFjwHjWzS Clh4Dk0qm5Y+m/T2AdX28a8chNDDz/Do9wAOCVelFSTXw121E0c3Y7kzoz38ZbL701QMCvQwwm 2tK/kIe44BChQ+R1usdI+qKae/6pltAcnFTT1qJYBnZRLtaGCd/I4EPoaXPupjtThXNApQFNan BuA= From: Thomas Schwinge To: Julian Brown CC: Jakub Jelinek , Subject: Re: [PATCH 1/2] [OpenACC] Refuse update/copyout for blocks with attached pointers In-Reply-To: <067e77d09132cbd32cc3f32c5af525f8edc2f53a.1592826181.git.julian@codesourcery.com> References: <067e77d09132cbd32cc3f32c5af525f8edc2f53a.1592826181.git.julian@codesourcery.com> User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.3 (x86_64-pc-linux-gnu) Date: Fri, 24 Jul 2020 15:37:08 +0200 Message-ID: <875zad6nln.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable 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, 24 Jul 2020 13:37:25 -0000 Hi Julian! Quoting your parent email: On 2020-06-22T05:14:42-0700, Julian Brown wrote: > Investigating PR95590, I realised that we can do better at diagnosing > some potentially troublesome usage of OpenACC "attach" behaviour, namely > updating blocks with attached pointers. Updating either the host copy > or device copy of such a block is problematic -- for a host update, > the host may get a clobbered (device) version of a host pointer in its > local version of the block (e.g. struct). A device update may clobber > an attached device pointer with a host pointer. ACK. > The spec text (OpenACC 3.0, "2.6.8. Attachment Counter") covering this > case is: > > "Pointer members of structs, classes, or derived types in device > or host memory can be overwritten due to update directives or API > routines. It is the user=E2=80=99s responsibility to ensure that the p= ointers > have the appropriate values before or after the data movement in > either direction. The behavior of the program is undefined if any > of the pointer members are attached when an update of a composite > variable is performed." > > The first patch in this series addresses that paragraph by making > such updates (as well as copyouts, similarly) be runtime errors. Hmm. But why do you say "addresses [...] by making [...] be runtime errors" if the specification text *explicitly* states ("It is the user's responsibility") that doing such things invokes undefined behavior, and thus a user must not do that. (Here, the undefined behavior is: copying of host vs. device pointers -- I wouldn't assume (user), respectively imply (implementor) anything worse?) It's of couse good if we can (without much overhead) be helpful to the user (your proposed runtime error), but I want to make sure that I'm correctly understanding your rationale here. On 2020-06-22T05:14:43-0700, Julian Brown wrote: > As mentioned in the parent email, this patch adds diagnostics for > probably-broken code that updates (host/device) or copies-out blocks > that still have attached pointers. Several new tests have been added. > > OK? I so far haven't managed to really convince myself that we want to incur this overhead here. (I suppose it's not too much overhead.) I may re-consider this still. I suppose we can put this onto the backburner -- nothing else functionally depends on this? Assuming this checking does get installed (and enabled by default), I had the idea that we may (rather easily?) add a flag variable (ICV; initialized from an environment variable) to guard this checking behavior? I suppose we may now have a few libgomp testcases that actually do use 'acc_update_self' etc. to read out pointer values from visible device copies, and verify these, which wouldn't work any longer with that checking enabled. Such tests could then 'dg-set-target-env-var "GOMP_ATTACH_CHECKING" "0"' (better name is desirable), and have one variant with and one variant without the checking. Gr=C3=BC=C3=9Fe Thomas > libgomp/ > * oacc-mem.c (update_dev_host): Raise error on update of block with > attached pointers. > (goacc_exit_data_internal): Raise error on copyout of block with > attached pointers. > * target.c (gomp_unmap_vars_internal): Likewise. > * testsuite/libgomp.oacc-c-c++-common/copyback-attached-dynamic-1.c= : > New test. > * testsuite/libgomp.oacc-c-c++-common/delete-attached-dynamic-1.c: > New test. > * testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-= 1.c: > New test. > * testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-= 2.c: > New test. > * testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-= 3.c: > New test. > * testsuite/libgomp.oacc-c-c++-common/delete-attached-structural-1.= c: > New test. > * testsuite/libgomp.oacc-c-c++-common/update-attached.c: New test. > * testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90: Updat= e > for new diagnostic. > --- > libgomp/oacc-mem.c | 42 ++++++++++++++++--- > libgomp/target.c | 27 +++++++++--- > .../copyback-attached-dynamic-1.c | 31 ++++++++++++++ > .../copyback-attached-structural-1.c | 30 +++++++++++++ > .../copyback-attached-structural-2.c | 31 ++++++++++++++ > .../copyback-attached-structural-3.c | 26 ++++++++++++ > .../delete-attached-dynamic-1.c | 26 ++++++++++++ > .../delete-attached-structural-1.c | 25 +++++++++++ > .../delete-attached-structural-2.c | 26 ++++++++++++ > .../update-attached-1.c | 33 +++++++++++++++ > .../deep-copy-6-no_finalize.F90 | 6 +-- > 11 files changed, 290 insertions(+), 13 deletions(-) > create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-= attached-dynamic-1.c > create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-= attached-structural-1.c > create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-= attached-structural-2.c > create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-= attached-structural-3.c > create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/delete-at= tached-dynamic-1.c > create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/delete-at= tached-structural-1.c > create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/delete-at= tached-structural-2.c > create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/update-at= tached-1.c > > diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c > index 1816b06bf2d..cf054f14b12 100644 > --- a/libgomp/oacc-mem.c > +++ b/libgomp/oacc-mem.c > @@ -865,6 +865,23 @@ update_dev_host (int is_dev, void *h, size_t s, int = async) > gomp_fatal ("[%p,%d] is not mapped", h, (int)s); > } > > + if (n->aux && n->aux->attach_count) > + { > + size_t nptrs =3D (n->host_end - n->host_start + sizeof (void *) - = 1) > + / sizeof (void *); > + for (size_t i =3D 0; i < nptrs; i++) > + if (n->aux->attach_count[i] > 0) > + { > + gomp_mutex_unlock (&acc_dev->lock); > + if (is_dev) > + gomp_fatal ("[%p,+%d] device update would overwrite attached = " > + "pointers", h, (int) s); > + else > + gomp_fatal ("host update from block [%p,+%d] with attached " > + "pointers", h, (int) s); > + } > + } > + > d =3D (void *) (n->tgt->tgt_start + n->tgt_offset > + (uintptr_t) h - n->host_start); > > @@ -1329,11 +1346,26 @@ goacc_exit_data_internal (struct gomp_device_desc= r *acc_dev, size_t mapnum, > if (copyfrom > && n->refcount !=3D REFCOUNT_INFINITY > && (kind !=3D GOMP_MAP_FROM || n->refcount =3D=3D 0)) > - gomp_copy_dev2host (acc_dev, aq, (void *) cur_node.host_start= , > - (void *) (n->tgt->tgt_start + n->tgt_offs= et > - + cur_node.host_start > - - n->host_start), > - cur_node.host_end - cur_node.host_start); > + { > + if (n->aux && n->aux->attach_count) > + { > + size_t nptrs =3D (n->host_end - n->host_start > + + sizeof (void *) - 1) / sizeof (void *= ); > + for (size_t j =3D 0; j < nptrs; j++) > + if (n->aux->attach_count[j] > 0) > + { > + gomp_mutex_unlock (&acc_dev->lock); > + gomp_fatal ("copyout of block [%p,+%d] with " > + "attached pointers", hostaddrs[i], > + (int) size); > + } > + } > + gomp_copy_dev2host (acc_dev, aq, (void *) cur_node.host_sta= rt, > + (void *) (n->tgt->tgt_start + n->tgt_of= fset > + + cur_node.host_start > + - n->host_start), > + cur_node.host_end - cur_node.host_start= ); > + } > > if (n->refcount =3D=3D 0) > { > diff --git a/libgomp/target.c b/libgomp/target.c > index badc254a777..db6f56a8ff8 100644 > --- a/libgomp/target.c > +++ b/libgomp/target.c > @@ -1649,11 +1649,28 @@ gomp_unmap_vars_internal (struct target_mem_desc = *tgt, bool do_copyfrom, > > if ((do_unmap && do_copyfrom && tgt->list[i].copy_from) > || tgt->list[i].always_copy_from) > - gomp_copy_dev2host (devicep, aq, > - (void *) (k->host_start + tgt->list[i].offset), > - (void *) (k->tgt->tgt_start + k->tgt_offset > - + tgt->list[i].offset), > - tgt->list[i].length); > + { > + if (k->aux && k->aux->attach_count) > + { > + size_t nptrs =3D (k->host_end - k->host_start > + + sizeof (void *) - 1) / sizeof (void *); > + for (size_t j =3D 0; j < nptrs; j++) > + if (k->aux->attach_count[j] > 0) > + { > + gomp_mutex_unlock (&devicep->lock); > + gomp_fatal ("copyout of block [%p,+%d] with " > + "attached pointers", > + (void *) (k->host_start + tgt->list[i].offs= et), > + (int) (k->host_end - k->host_start)); > + } > + } > + gomp_copy_dev2host (devicep, aq, > + (void *) (k->host_start + tgt->list[i].offset= ), > + (void *) (k->tgt->tgt_start + k->tgt_offset > + + tgt->list[i].offset), > + tgt->list[i].length); > + } > + > if (do_unmap) > { > struct target_mem_desc *k_tgt =3D k->tgt; > diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attache= d-dynamic-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attach= ed-dynamic-1.c > new file mode 100644 > index 00000000000..bc4e297fa6f > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-dynam= ic-1.c > @@ -0,0 +1,31 @@ > +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=3D1" } } */ > + > +#include > + > +struct mystruct { > + int *arr; > +}; > + > +int > +main (int argc, char *argv[]) > +{ > + int localarray[1024]; > + struct mystruct s; > + s.arr =3D localarray; > + > + #pragma acc enter data copyin(s) > + > + #pragma acc data copy(s.arr[0:1024]) > + { > + /* This directive does one too many attachments: it should fail when= we try > + to do the copyout below. */ > + #pragma acc enter data attach(s.arr) > + /* { dg-output "copyout of block \\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\] w= ith attached pointers" } */ > + } > + > + #pragma acc exit data copyout(s) > + > + return 0; > +} > + > +/* { dg-shouldfail "" } */ > diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attache= d-structural-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-att= ached-structural-1.c > new file mode 100644 > index 00000000000..7846c8c717c > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-struc= tural-1.c > @@ -0,0 +1,30 @@ > +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=3D1" } } */ > + > +#include > + > +struct mystruct { > + int *arr; > +}; > + > +int > +main (int argc, char *argv[]) > +{ > + int localarray[1024]; > + struct mystruct s; > + s.arr =3D localarray; > + > + #pragma acc data copy(s) > + { > + #pragma acc data copy(s.arr[0:1024]) > + { > + /* This directive does one too many attachments: it should fail wh= en we try > + to do the copyout below. */ > + #pragma acc enter data attach(s.arr) > + /* { dg-output "copyout of block \\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\]= with attached pointers" } */ > + } > + } > + > + return 0; > +} > + > +/* { dg-shouldfail "" } */ > diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attache= d-structural-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-att= ached-structural-2.c > new file mode 100644 > index 00000000000..bffa06eb725 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-struc= tural-2.c > @@ -0,0 +1,31 @@ > +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=3D1" } } */ > + > +#include > + > +struct mystruct { > + int *arr; > +}; > + > +int > +main (int argc, char *argv[]) > +{ > + int localarray[1024]; > + struct mystruct s; > + s.arr =3D localarray; > + > + #pragma acc enter data copyin(localarray[0:1024]) > + > + #pragma acc data copy(s) > + { > + /* This directive does one too many attachments: it should fail when= we try > + to do the copyout below. */ > + #pragma acc enter data attach(s.arr) > + /* { dg-output "copyout of block \\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\] w= ith attached pointers" } */ > + } > + > + #pragma acc exit data delete(localarray[0:1024]) > + > + return 0; > +} > + > +/* { dg-shouldfail "" } */ > diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attache= d-structural-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-att= ached-structural-3.c > new file mode 100644 > index 00000000000..4b21677af09 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-struc= tural-3.c > @@ -0,0 +1,26 @@ > +#include > + > +struct mystruct { > + int *arr; > +}; > + > +int > +main (int argc, char *argv[]) > +{ > + int localarray[1024]; > + struct mystruct s; > + s.arr =3D localarray; > + > + #pragma acc enter data copyin(localarray[0:1024]) > + > + #pragma acc data copy(s) > + { > + /* Here the attach and detach balance: this should work. */ > + #pragma acc enter data attach(s.arr) > + #pragma acc exit data detach(s.arr) > + } > + > + #pragma acc exit data delete(localarray[0:1024]) > + > + return 0; > +} > diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-= dynamic-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-d= ynamic-1.c > new file mode 100644 > index 00000000000..e074d507fb2 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-dynamic= -1.c > @@ -0,0 +1,26 @@ > +#include > + > +struct mystruct { > + int *arr; > +}; > + > +int > +main (int argc, char *argv[]) > +{ > + int localarray[1024]; > + struct mystruct s; > + s.arr =3D localarray; > + > + #pragma acc enter data copyin(s) > + > + #pragma acc data copy(s.arr[0:1024]) > + { > + /* We delete 's' from the target below: this extra attachment is not > + dangerous and we do not raise an error. */ > + #pragma acc enter data attach(s.arr) > + } > + > + #pragma acc exit data delete(s) > + > + return 0; > +} > diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-= structural-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attache= d-structural-1.c > new file mode 100644 > index 00000000000..e675762ecd8 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-structu= ral-1.c > @@ -0,0 +1,25 @@ > +#include > + > +struct mystruct { > + int *arr; > +}; > + > +int > +main (int argc, char *argv[]) > +{ > + int localarray[1024]; > + struct mystruct s; > + s.arr =3D localarray; > + > + #pragma acc data copyin(s) > + { > + #pragma acc data copy(s.arr[0:1024]) > + { > + /* This directive does one too many attachments: it should fail wh= en we try > + to do the copyout below. */ > + #pragma acc enter data attach(s.arr) > + } > + } > + > + return 0; > +} > diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-= structural-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attache= d-structural-2.c > new file mode 100644 > index 00000000000..d2095255ad3 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-structu= ral-2.c > @@ -0,0 +1,26 @@ > +#include > + > +struct mystruct { > + int *arr; > +}; > + > +int > +main (int argc, char *argv[]) > +{ > + int localarray[1024]; > + struct mystruct s; > + s.arr =3D localarray; > + > + #pragma acc enter data copyin(localarray[0:1024]) > + > + #pragma acc data copyin(s) > + { > + /* We only try to copy in: the extra attachment we're left over with= is not > + harmful and we don't raise an error. */ > + #pragma acc enter data attach(s.arr) > + } > + > + #pragma acc exit data delete(localarray[0:1024]) > + > + return 0; > +} > diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/update-attached-= 1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/update-attached-1.c > new file mode 100644 > index 00000000000..9f60bfa56f4 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/update-attached-1.c > @@ -0,0 +1,33 @@ > +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=3D1" } } */ > + > +#include > + > +struct mystruct { > + int *arr; > +}; > + > +int > +main (int argc, char *argv[]) > +{ > + int localarray[1024]; > + int localarray2[1024]; > + struct mystruct s; > + s.arr =3D localarray; > + > + #pragma acc enter data copyin(s) > + > + #pragma acc data copy(s.arr[0:1024]) > + { > + s.arr =3D localarray2; > + /* This update is dangerous because we have attached pointers: raise= an > + error. */ > + #pragma acc update device(s) > + /* { dg-output "\\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\] device update woul= d overwrite attached pointers" } */ > + } > + > + #pragma acc exit data delete(s) > + > + return 0; > +} > + > +/* { dg-shouldfail "" } */ > diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finali= ze.F90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90 > index ad8da71d7c9..355a381b625 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 > @@ -8,7 +8,7 @@ > ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" } > ! { dg-output ".*CheCKpOInT2(\n|\r\n|\r)" } > > -! Without the finalize, we do not detach properly so the host sees a dev= ice > -! pointer, and fails with this STOP code. > -! { dg-output "STOP 7(\n|\r\n|\r)+" } > +! Without the finalize, we do not detach properly and raise an error on = attempting > +! the copyout. > +! { dg-output ".*copyout of block \\\[0x\[0-9a-f\]+,\\+\[0-9\]+\\\] with= attached pointers(\n|\r\n|\r)+" } > ! { dg-shouldfail "" } ----------------- 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