From patchwork Mon Jun 22 12:14:43 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1314335 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 49r7dx3832z9sQx for ; Mon, 22 Jun 2020 22:15:17 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 946B53840C0F; Mon, 22 Jun 2020 12:15:09 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 76B3E3840C0A for ; Mon, 22 Jun 2020 12:15:05 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 76B3E3840C0A 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: opIBKNAlncgWPrIv8O81MKTH5vMJfIJmcocsvGKy8abBjKPuZSYtRJB0IOPdi5oFJS9q1HwLx9 L8Z7xS6LYosI7BcfJc+TZfd3rMIZrvnI2aIkpdCVpoo4EAkBnv8cLsTSK9IVV34UOSBw314oWq oHGwxUsVQ2t0c1Ha4mDQm7GBMpun9S5EHwBY4tcNrYyGvY3cKPcKaLHBmLSWUlg8z+/xByQJDh hXQD2KAtlGdLAGBNuy4L+KZy40qsziQvuUDZeMedw/pCpfOzunUlVL/SXEZR9ITTC3Yg9dODMb hdw= X-IronPort-AV: E=Sophos;i="5.75,266,1589270400"; d="scan'208";a="50176515" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 22 Jun 2020 04:15:05 -0800 IronPort-SDR: SsUU0EGjoUBwEI5H/dv7spv8RI++K49pzkSiAKCG3FsNNLC67cHn251x4lcgPCIyY4qZ9C29+W IOudj5g1eWa/wGdSBFqk+juR03LgpZw8iL9O6Cof1ZKZR6Qfe6nGhJGnczK7O1jZhlclgt7v8i +kaKu1y3+yk+9+QDO88T96IYU+9PRALxoj9FmGOTouoejR0jhIY6wvqT7+5q1WzQSz7BrBZYh2 JQ1cjiP8GxOJPBBxI8ASUQCUpp4HV4QaTRc5TF0GWL43clr7vtmaaej7buHSdPyq9F0b/lwJRO J7s= From: Julian Brown To: Subject: [PATCH 1/2] [OpenACC] Refuse update/copyout for blocks with attached pointers Date: Mon, 22 Jun 2020 05:14:43 -0700 Message-ID: <067e77d09132cbd32cc3f32c5af525f8edc2f53a.1592826181.git.julian@codesourcery.com> X-Mailer: git-send-email 2.23.0 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-12.7 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: , Cc: Jakub Jelinek , Thomas Schwinge Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" 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? Julian ChangeLog 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: Update 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-attached-dynamic-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-structural-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-structural-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/update-attached-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 = (n->host_end - n->host_start + sizeof (void *) - 1) + / sizeof (void *); + for (size_t i = 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 = (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_descr *acc_dev, size_t mapnum, if (copyfrom && n->refcount != REFCOUNT_INFINITY && (kind != GOMP_MAP_FROM || n->refcount == 0)) - gomp_copy_dev2host (acc_dev, aq, (void *) cur_node.host_start, - (void *) (n->tgt->tgt_start + n->tgt_offset - + 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 = (n->host_end - n->host_start + + sizeof (void *) - 1) / sizeof (void *); + for (size_t j = 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_start, + (void *) (n->tgt->tgt_start + n->tgt_offset + + cur_node.host_start + - n->host_start), + cur_node.host_end - cur_node.host_start); + } if (n->refcount == 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 = (k->host_end - k->host_start + + sizeof (void *) - 1) / sizeof (void *); + for (size_t j = 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].offset), + (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 = k->tgt; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-dynamic-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-dynamic-1.c new file mode 100644 index 00000000000..bc4e297fa6f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-dynamic-1.c @@ -0,0 +1,31 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include + +struct mystruct { + int *arr; +}; + +int +main (int argc, char *argv[]) +{ + int localarray[1024]; + struct mystruct s; + s.arr = 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\]+\\\] with attached pointers" } */ + } + + #pragma acc exit data copyout(s) + + return 0; +} + +/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-1.c new file mode 100644 index 00000000000..7846c8c717c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-1.c @@ -0,0 +1,30 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include + +struct mystruct { + int *arr; +}; + +int +main (int argc, char *argv[]) +{ + int localarray[1024]; + struct mystruct s; + s.arr = localarray; + + #pragma acc data copy(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\]+\\\] with attached pointers" } */ + } + } + + return 0; +} + +/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-2.c new file mode 100644 index 00000000000..bffa06eb725 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-2.c @@ -0,0 +1,31 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include + +struct mystruct { + int *arr; +}; + +int +main (int argc, char *argv[]) +{ + int localarray[1024]; + struct mystruct s; + s.arr = 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\]+\\\] with 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-attached-structural-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-3.c new file mode 100644 index 00000000000..4b21677af09 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyback-attached-structural-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 = 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-dynamic-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 = 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-attached-structural-1.c new file mode 100644 index 00000000000..e675762ecd8 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-structural-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 = localarray; + + #pragma acc 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) + } + } + + 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-attached-structural-2.c new file mode 100644 index 00000000000..d2095255ad3 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/delete-attached-structural-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 = 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=1" } } */ + +#include + +struct mystruct { + int *arr; +}; + +int +main (int argc, char *argv[]) +{ + int localarray[1024]; + int localarray2[1024]; + struct mystruct s; + s.arr = localarray; + + #pragma acc enter data copyin(s) + + #pragma acc data copy(s.arr[0:1024]) + { + s.arr = 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 would 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_finalize.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 device -! 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 "" } From patchwork Mon Jun 22 12:14:44 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1314336 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 49r7f01q4wz9sQx for ; Mon, 22 Jun 2020 22:15:20 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 04E313840C26; Mon, 22 Jun 2020 12:15:10 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id D64673851C03 for ; Mon, 22 Jun 2020 12:15:07 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org D64673851C03 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: 7nxmHviyQuPwo2l7qJKR3DGsn8P7J9vjVpKC7lYdi9O2IoGpgKLv1PH8RGd+5wMyCTLp3+S9h0 TqD6GjlDnfpKeQdELkYMLaqHaR332dqB40A3uArCeIt9HMeHyrLGAX4V6PyIS3Lm7pSpilZOsX ED7wZr3LvKOsJscGojVWqTwu74oL6wLKmjsSP6kUiHK9aUetcKDBXyBIgmR/4Ea3j4uhTCON1c 33UGPKW1Dm5iQKMOpDQF6J3xMq9Wug/Ox/8NfgI0YIOJiHUtfQ3oW4ht8NJ2XTfSDF+B9lpFKh QzE= X-IronPort-AV: E=Sophos;i="5.75,266,1589270400"; d="scan'208";a="50176518" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 22 Jun 2020 04:15:07 -0800 IronPort-SDR: 8W8pboPykrgjgoMLqMaXx4amFbuYd3dYyyLfD6FKC7LojQPa3y8zLrSJ3OvhRHckk6WcO3bNX/ tt6B1pNVJyUEM5DR9gDj0eseoLw16oMtO0oIyQlVo2t7ka8gzQGJGW0rCKTA4eXLa/sZ6zc39e G5vRiAG6pEtMTzuZGXQG+AGbZT7yyfink2boeAugeHqsm/V92Ct5qDr3+c8OwS5MLYop/zqW2H G8ZB+w074YY2LKx1F56v1elDJpW8n8j8kHz6CJsdseFkykCix531JrBRNr+uFKlpl77fQRGC9n ydE= From: Julian Brown To: Subject: [PATCH 2/2] [OpenACC] Detect pointer updates for attach operations (PR95590) Date: Mon, 22 Jun 2020 05:14:44 -0700 Message-ID: X-Mailer: git-send-email 2.23.0 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-12.7 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: , Cc: Jakub Jelinek , Thomas Schwinge Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" As mentioned in the parent email, this is a fix for PR95590 that detects updates of attached pointers in blocks, and rewrites the attached pointer and resets its attachment counter appropriately. I am however not entirely sure this is desirable or required by the spec: points against are: - To avoid expensive copies from the device to the host and/or "wrong way" device-to-host splay tree lookups, it requires keeping an extra shadow copy of mapped blocks on the host in order to detect if a host pointer with attachments in the block has been changed between attach operations. We incur this overhead unconditionally if attach/detach are in use for what's not likely to be a common use case (it's slightly tricky to write a test case to exercise the behaviour, even -- Thomas's unmodified original for the PR raises an error after the previous patch in this series). - From a user perspective, I think it's going to be quite easy to get confused wrt. the hidden attachment counter state, with this kind of reset-on-host-pointer-modification behaviour. Mind you, silently *not* doing the update is likewise going to be confusing (the stale device pointer would be updated at present). Maybe this should be detected as an error instead? - The text in "2.6.8. Attachment Counter" *might* contribute to the argument that this kind of pointer-update detection is not required. Anyway, thoughts, or OK for mainline? Thanks, Julian ChangeLog PR libgomp/95590 libgomp/ * target.c (gomp_attach_pointer): Initialise shadow copy of block with attached pointers, and use to detect modifications of those pointers. * testsuite/libgomp.oacc-c-c++-common/attach-ptr-change-1.c: New test. --- libgomp/target.c | 29 +++++++- .../attach-ptr-change-1.c | 74 +++++++++++++++++++ 2 files changed, 100 insertions(+), 3 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/attach-ptr-change-1.c diff --git a/libgomp/target.c b/libgomp/target.c index db6f56a8ff8..076cc2bbbcb 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -691,6 +691,8 @@ gomp_attach_pointer (struct gomp_device_descr *devicep, { struct splay_tree_key_s s; size_t size, idx; + char *shadow_block; + size_t shadow_size = n->host_end - n->host_start; if (n == NULL) { @@ -707,9 +709,31 @@ gomp_attach_pointer (struct gomp_device_descr *devicep, if (!n->aux) n->aux = gomp_malloc_cleared (sizeof (struct splay_tree_aux)); + bool first = false; + if (!n->aux->attach_count) - n->aux->attach_count - = gomp_malloc_cleared (sizeof (*n->aux->attach_count) * size); + { + n->aux->attach_count + = gomp_malloc_cleared (sizeof (*n->aux->attach_count) * size + + shadow_size); + first = true; + } + + shadow_block = ((char *) n->aux->attach_count) + + sizeof (*n->aux->attach_count) * size; + + if (first) + memcpy (shadow_block, (const void *) n->host_start, shadow_size); + + uintptr_t target = (uintptr_t) *(void **) attach_to; + uintptr_t shadow_target + = (uintptr_t) *(void **) (shadow_block + attach_to - n->host_start); + if (target != shadow_target) + { + n->aux->attach_count[idx] = 0; + memcpy ((char *) shadow_block + attach_to - n->host_start, + (const void *) target, sizeof (void *)); + } if (n->aux->attach_count[idx] < UINTPTR_MAX) n->aux->attach_count[idx]++; @@ -723,7 +747,6 @@ gomp_attach_pointer (struct gomp_device_descr *devicep, { uintptr_t devptr = n->tgt->tgt_start + n->tgt_offset + attach_to - n->host_start; - uintptr_t target = (uintptr_t) *(void **) attach_to; splay_tree_key tn; uintptr_t data; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/attach-ptr-change-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/attach-ptr-change-1.c new file mode 100644 index 00000000000..d4d84fdb092 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/attach-ptr-change-1.c @@ -0,0 +1,74 @@ +#include +#include +#include + +struct str { + unsigned char *c; +}; + +int main() +{ + const int size_1 = sizeof (void *); + unsigned char *data_1 = (unsigned char *) malloc(sizeof (void *)); + assert(data_1); + void *data_1_d = acc_create(data_1, size_1); + assert(data_1_d); + assert(acc_is_present(data_1, size_1)); + + const int size_2 = sizeof (void *); + unsigned char *data_2 = (unsigned char *) malloc(size_2); + assert(data_2); + void *data_2_d = acc_create(data_2, size_2); + assert(data_2_d); + assert(acc_is_present(data_2, size_2)); + + struct str data_work; + data_work.c = data_1; + + acc_copyin(&data_work, sizeof data_work); + assert(acc_is_present(&data_work, sizeof data_work)); + assert(data_work.c == data_1); + + /* No attach has taken place so far. We can still do a self-update. */ + acc_update_self(&data_work, sizeof data_work); + assert(data_work.c == data_1); + + data_1[0] = 'a'; + data_2[0] = 'b'; + + acc_update_device (data_1, size_1); + acc_update_device (data_2, size_2); + + acc_attach((void **) &data_work.c); + #pragma acc serial /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ + { + data_work.c[0] = 'c'; + } + + acc_update_self (data_1, size_1); + acc_update_self (data_2, size_2); + + assert (data_1[0] == 'c'); + assert (data_2[0] == 'b'); + + data_1[0] = 'a'; + data_2[0] = 'b'; + + acc_update_device (data_1, size_1); + acc_update_device (data_2, size_2); + + data_work.c = data_2; + acc_attach((void **) &data_work.c); + #pragma acc serial /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ + { + data_work.c[0] = 'd'; + } + + acc_update_self (data_1, size_1); + acc_update_self (data_2, size_2); + + assert (data_1[0] == 'a'); + assert (data_2[0] == 'd'); + + return 0; +}