From patchwork Sat Nov 9 01:04:21 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1192348 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=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-512869-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="y60WRJC9"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 478zSV1yp6z9sPF for ; Sat, 9 Nov 2019 12:04:42 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:subject:message-id:mime-version:content-type; q=dns; s= default; b=v79CRo7oAWwqD/Pi3pyCiE3HnraokeP7jg+R8ABHBEIrms6l0zOeT JerbJ6oH6mDN7h/B4HThWw0Xd6AeqIy9eTJFjhU+4gq7XVyByKaVGHY17SWMgaYF mFaKQRjRKMwvCtecNQaTrflh+p+7KtdM752IMoOm020Kup3ns2W+e8= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:subject:message-id:mime-version:content-type; s= default; bh=By1ePttCtps7ecIaTFFt+qFcfJg=; b=y60WRJC96CiVh2Ai045g a30csozoVEXTkMUu+eu17Z5xrB1q+Z3g+MpFUxi2/5LiQZiylStpgR+vefAhVT59 jQzkXY+VuQ8uXj2G7IGxsctXLeluAwTktwYHFPoX7Rq9ZQvGy2ZZP2kCDdm+/CDa YzedtYs5y8EG2D+J6nVOL7I= Received: (qmail 78019 invoked by alias); 9 Nov 2019 01:04:35 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 78010 invoked by uid 89); 9 Nov 2019 01:04:34 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-15.9 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, KAM_SHORT autolearn=ham version=3.3.1 spammy=5357, subarrays X-HELO: esa4.mentor.iphmx.com Received: from esa4.mentor.iphmx.com (HELO esa4.mentor.iphmx.com) (68.232.137.252) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sat, 09 Nov 2019 01:04:32 +0000 IronPort-SDR: oWX57mXmiWDCah9zBxSiTMTKp7dum0AVuCes8y7UDpUih4y8BdGL9D9DEYr7vphEF/H0iHlMHV KjwGQsXWhVkPrCCqS05lT1O5O5NNeIR4N5u3dCMTjuSp1UvhYCMT8xJJzDdlpdN7Qtr6NySAH8 FIDodHGJ/jZhAknCACloY+4qGRAtoGqhel5SWDKvOxfRrzgqpTX7IkYTe3ENPNzON9yQgg89NN CpCmJM7HzZ/xW6ZJ/cyEFMhuBzZeelSMX0td1rojsRKoBaDzAoMWaBWD0mPRAZLNCda4VKPGBx 6eQ= Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 08 Nov 2019 17:04:31 -0800 IronPort-SDR: hr3Mab5+MyXCRoSpMsNDsiXVufjG47wSK2h3E9PcFCYNB1tAx/TSSuZQ37m3ElnARX4xKu1oMa 76U91Q6RSKcgxun+27uYD3DGn+rhcu1CCfdTi4y4xV50I7Jt1FAfUSkf0fJGcWZv1sYMFkp1DO Qzp/VApuGvMU9K3u+rAAEnknQY5/jHsZuqort80S6hDFxdH/qeN82to7v3yV+3JmTGh1ikNGBd 5kah1aWB8KwEboWeuEdPnese4I5RszyWJJxYd2+ykNo0FgUidVct8LlPRJTWfMsZHgDDNmfsu1 1Sg= Date: Sat, 9 Nov 2019 01:04:21 +0000 From: Julian Brown To: Thomas Schwinge , Jakub Jelinek , Subject: [PATCH] OpenACC "present" subarrays: runtime API return value and unmapping fixes Message-ID: <20191109010421.5b8b689c@squid.athome> MIME-Version: 1.0 X-IsSubscribed: yes Hi, This patch fixes an issue I noticed when investigating an answer for Thomas's question about device pointer return values in: https://gcc.gnu.org/ml/gcc-patches/2019-10/msg02260.html It looks to me like the return value for the present case is wrong in the existing code: in case of a acc_pcopyin or similar call that refers to a subarray of a larger block already mapped on the target, the device pointer return value will be the start of the larger block, not of the subarray being copied. The attached patch corrects this issue, and also relaxes a restriction on acc_delete, acc_copyout (etc.) to allow them to unmap/copyout subarrays of a larger block already present on the target. There's no particular reason to disallow that, as far as I can tell. This is necessary to allow the new tests included with this patch to pass, and a couple of existing "shouldfail" tests no longer fail, and have been adjusted accordingly. It's still an error to try to copy data beyond the bounds of a mapped block, and other existing tests cover those cases. The calculation for the return value for the non-present case of present_create_copy has also been adjusted in anticipation of a new version of the above-linked patch. Tested with offloading to nvptx. OK for trunk? Julian ChangeLog libgomp/ * oacc-mem.c (present_create_copy): Fix device pointer return value in case of "present" subarray. Use tgt->tgt_start instead of tgt->to_free in non-present/create case. (delete_copyout): Change error condition to detect only copies outside of mapped block. Adjust error message accordingly. * testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/lib-20.c: Adjust expected error message. * testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-22.c: Allow test to pass now. * testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise. commit 00607b06c8e506b0f0744a230856e1e8776633c3 Author: Julian Brown Date: Thu Nov 7 14:24:49 2019 -0800 OpenACC "present" subarrays: runtime API return value and unmapping fixes libgomp/ * oacc-mem.c (present_create_copy): Fix device pointer return value in case of "present" subarray. Use tgt->tgt_start instead of tgt->to_free in non-present/create case. (delete_copyout): Change error condition to fail only on copies outside of mapped block. Adjust error message accordingly. * testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/lib-20.c: Adjust expected error message. * testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-22.c: Allow test to pass now. * testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise. diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 2f271009fb8..0a41f11210c 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -535,7 +535,7 @@ present_create_copy (unsigned f, void *h, size_t s, int async) if (n) { /* Present. */ - d = (void *) (n->tgt->tgt_start + n->tgt_offset); + d = (void *) (n->tgt->tgt_start + n->tgt_offset + h - n->host_start); if (!(f & FLAG_PRESENT)) { @@ -584,7 +584,7 @@ present_create_copy (unsigned f, void *h, size_t s, int async) gomp_mutex_lock (&acc_dev->lock); - d = tgt->to_free; + d = (void *) tgt->tgt_start; tgt->prev = acc_dev->openacc.data_environ; acc_dev->openacc.data_environ = tgt; @@ -669,7 +669,6 @@ acc_pcopyin (void *h, size_t s) static void delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname) { - size_t host_size; splay_tree_key n; void *d; struct goacc_thread *thr = goacc_thread (); @@ -703,13 +702,12 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname) d = (void *) (n->tgt->tgt_start + n->tgt_offset + (uintptr_t) h - n->host_start); - host_size = n->host_end - n->host_start; - - if (n->host_start != (uintptr_t) h || host_size != s) + if ((uintptr_t) h < n->host_start || (uintptr_t) h + s > n->host_end) { + size_t host_size = n->host_end - n->host_start; gomp_mutex_unlock (&acc_dev->lock); - gomp_fatal ("[%p,%d] surrounds2 [%p,+%d]", - (void *) n->host_start, (int) host_size, (void *) h, (int) s); + gomp_fatal ("[%p,+%d] outside mapped block [%p,+%d]", + (void *) h, (int) s, (void *) n->host_start, (int) host_size); } if (n->refcount == REFCOUNT_INFINITY) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c new file mode 100644 index 00000000000..bee0b10ca7b --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c @@ -0,0 +1,28 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include +#include + +int main (int argc, char* argv[]) +{ + char* myblock = malloc (1024); + int i; + void *dst; + for (i = 0; i < 1024; i++) + myblock[i] = i; + dst = acc_copyin (myblock, 1024); + for (i = 0; i < 1024; i += 256) + { + void *partdst = acc_pcopyin (&myblock[i], 256); + assert ((uintptr_t) partdst == (uintptr_t) dst + i); + } + for (i = 0; i < 1024; i += 256) + acc_delete (&myblock[i], 256); + assert (acc_is_present (myblock, 1024)); + acc_delete (myblock, 1024); + assert (!acc_is_present (myblock, 1024)); + free (myblock); + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c new file mode 100644 index 00000000000..d35ab5c4b71 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c @@ -0,0 +1,35 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include +#include + +int main (int argc, char* argv[]) +{ + char* block1 = malloc (1024); + char *block2 = malloc (1024); + char *block3 = malloc (1024); + int i; + void *dst; + for (i = 0; i < 1024; i++) + block1[i] = block2[i] = block3[i] = i; + #pragma acc data copyin(block1[0:1024]) copyin(block2[0:1024]) \ + copyin(block3[0:1024]) + { + dst = acc_deviceptr (block2); + for (i = 0; i < 1024; i += 256) + { + void *partdst = acc_pcopyin (&block2[i], 256); + assert ((uintptr_t) partdst == (uintptr_t) dst + i); + } + } + assert (acc_is_present (block2, 1024)); + for (i = 0; i < 1024; i += 256) + acc_delete (&block2[i], 256); + assert (!acc_is_present (block2, 1024)); + free (block1); + free (block2); + free (block3); + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c index 25ceb3a26af..10d3cbc5cc6 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c @@ -31,5 +31,5 @@ main (int argc, char **argv) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,256\\\] surrounds2 \\\[\[0-9a-fA-FxX\]+,\\\+257\\\]" } */ +/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+257\\\] outside mapped block \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */ /* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c index 65ff440a528..cb32bbcb652 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c @@ -31,5 +31,3 @@ main (int argc, char **argv) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,256\\\] surrounds2 \\\[\[0-9a-fA-FxX\]+,\\\+255\\\]" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c index fd4dc5971a1..b1f3e71f278 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c @@ -41,5 +41,5 @@ main (int argc, char **argv) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,256\\\] surrounds2 \\\[\[0-9a-fA-FxX\]+,\\\+512\\\]" } */ +/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+512\\\] outside mapped block \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */ /* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c index 9bc9ecc1068..d0e5ffb0691 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c @@ -28,5 +28,3 @@ main (int argc, char **argv) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,256\\\] surrounds2 \\\[\[0-9a-fA-FxX\]+,\\\+254\\\]" } */ -/* { dg-shouldfail "" } */