From patchwork Mon Dec 9 23:03:04 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1206784 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-515556-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="wl8mzkm9"; 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 47WzJL386jz9sNH for ; Tue, 10 Dec 2019 10:03:30 +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:from :to:cc:subject:date:message-id:mime-version:content-type; q=dns; s=default; b=p8LoC7FUGLU/jnXsdUg7b3+HqCHZiWTXokpKs3kE/viZKBNlCN ff6XpxeKogBZkgbMejYswVOtFhUGgY/gHP2DcAQKtjeeiiBieMJKswAOh2Os/IH+ ouFJ2RFo9DkxxsI8tIzdyMpjCdg6r7PBNd5kyYuBWqBLprjzwT82pM5Pk= 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:from :to:cc:subject:date:message-id:mime-version:content-type; s= default; bh=QrRvnWq6HJd9lop80400s+Otu00=; b=wl8mzkm9LOzcew4qpYh7 MFq3dwMntwATQDFCHoO/iAzJ3WLOxj8lbSmwCYoInr6YEdR+Oltdt0dfUveQOxj/ shvoi7bI2d0CbsDckKvR8v+VARbMJekyO7PTkpcJ0Z3ORW9CRb4llbopwMo2mCeU GQrIzz9o8A2JSZZV4TsO82s= Received: (qmail 79727 invoked by alias); 9 Dec 2019 23:03:22 -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 79714 invoked by uid 89); 9 Dec 2019 23:03:21 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-17.5 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SUBJ_OBFU_PUNCT_FEW, SUBJ_OBFU_PUNCT_MANY autolearn=ham version=3.3.1 spammy= X-HELO: esa2.mentor.iphmx.com Received: from esa2.mentor.iphmx.com (HELO esa2.mentor.iphmx.com) (68.232.141.98) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Mon, 09 Dec 2019 23:03:16 +0000 IronPort-SDR: CvtM6UXG/+R2Pbut5oPTVGi9HYmnHebwnCfTZ2HOZ5LX2usK1E8RWhSTHzL9dCXvGz/Fks0Fa9 XmBlua4W4nHhJ/cIrIZJNi/4vVNO1mGD+isLUTxzN7tyRqhvtFaBI+qJ9PwE/wMZ4iHVDWV0ZV 8mAZGBWkJ1Qrs1PnkRzcph4DUE62t4BTQgwVTWRz2/4s5EAXI0nHQPQEFDh7DTmRYZexCgTe5R 5Yq3Mie6Q5bTrn22eipkJW8ly41/aW24ddlA3GHDjGStzHTTTaI4KK6AB3Uf/7rRrXUzmr1EZX S2E= Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 09 Dec 2019 15:03:14 -0800 IronPort-SDR: 2Fjjo1TyPFUsT+gx83WEvKZ0VckJifbzSSXsKLnABlM2dvqy+IRGBtLv5hG3cpd4CCNK8HOvNz qZQ7nCC1fMBT2jMnLWJvHzO8FAlr0twE8f47wQKTflXdwkzAaHIkKD4ifYzJMIogTFSdQea2v+ d7bP6JDO5GJDRqD2g+Dlr0MoCfQ6fjUe8gd6M46fBNoY187t0kUp35Iyr8+JqJJxsPqPKO2DDg zej1wconNFDg331mT7ruy3Rn8qkwJ+vj6NB5S43x2yw4cZ7/54BQS9SEk8xXXauI59E+nFTuJo 3Sw= From: Thomas Schwinge To: CC: Jakub Jelinek , Julian Brown Subject: [PR92503] [OpenACC] Don't silently 'acc_unmap_data' in 'acc_free' User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.1 (x86_64-pc-linux-gnu) Date: Tue, 10 Dec 2019 00:03:04 +0100 Message-ID: <871rtd2jbr.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Hi! See attached "[PR92503] [OpenACC] Don't silently 'acc_unmap_data' in 'acc_free'", committed to trunk in r279146. As mentioned in PR92503, further work can be done later on, incrementally, to avoid "expensive device-to-host-address lookup": possibly "we might actually keep such additional/expensive sanity-checking, but guard it by an environment variable". Grüße Thomas From 03383a93c7318009ddd0e8d77b1a950c4b2b8f5a Mon Sep 17 00:00:00 2001 From: tschwinge Date: Mon, 9 Dec 2019 22:52:47 +0000 Subject: [PATCH] [PR92503] [OpenACC] Don't silently 'acc_unmap_data' in 'acc_free' libgomp/ PR libgomp/92503 * oacc-mem.c (acc_free): Error out instead of 'acc_unmap_data'. * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c: New file. * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Adjust. * testsuite/libgomp.oacc-c-c++-common/context-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/context-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/context-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/context-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-13.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-91.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@279146 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog | 25 +++++++++++++++ libgomp/oacc-mem.c | 17 +++++----- .../acc_free-pr92503-1.c | 28 ++++++++++++++++ .../acc_free-pr92503-2.c | 27 ++++++++++++++++ .../acc_free-pr92503-3-2.c | 28 ++++++++++++++++ .../acc_free-pr92503-3.c | 28 ++++++++++++++++ .../acc_free-pr92503-4-2.c | 31 ++++++++++++++++++ .../acc_free-pr92503-4.c | 32 +++++++++++++++++++ .../libgomp.oacc-c-c++-common/clauses-1.c | 12 +++++-- .../libgomp.oacc-c-c++-common/context-1.c | 6 ++-- .../libgomp.oacc-c-c++-common/context-2.c | 6 ++-- .../libgomp.oacc-c-c++-common/context-3.c | 6 ++-- .../libgomp.oacc-c-c++-common/context-4.c | 6 ++-- .../libgomp.oacc-c-c++-common/lib-13.c | 2 +- .../libgomp.oacc-c-c++-common/lib-14.c | 2 +- .../libgomp.oacc-c-c++-common/lib-18.c | 2 +- .../libgomp.oacc-c-c++-common/lib-91.c | 2 ++ .../libgomp.oacc-c-c++-common/nested-1.c | 12 +++++-- 18 files changed, 242 insertions(+), 30 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 7606f17825d..62092a2d765 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,5 +1,30 @@ 2019-12-09 Thomas Schwinge + PR libgomp/92503 + * oacc-mem.c (acc_free): Error out instead of 'acc_unmap_data'. + * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c: New + file. + * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c: + Likewise. + * testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Adjust. + * testsuite/libgomp.oacc-c-c++-common/context-1.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/context-2.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/context-3.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/context-4.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-13.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/lib-91.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise. + PR libgomp/92840 * oacc-mem.c (acc_map_data): Clarify reference counting behavior. (acc_unmap_data): Add error case for 'REFCOUNT_INFINITY'. diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 480b9fbb71b..81ebddf7580 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -121,9 +121,6 @@ acc_malloc (size_t s) return res; } -/* OpenACC 2.0a (3.2.16) doesn't specify what to do in the event - the device address is mapped. We choose to check if it mapped, - and if it is, to unmap it. */ void acc_free (void *d) { @@ -152,13 +149,15 @@ acc_free (void *d) (unless you got that null from acc_malloc). */ if ((k = lookup_dev (acc_dev->openacc.data_environ, d, 1))) { - void *offset; - - offset = d - k->tgt->tgt_start + k->tgt_offset; - + void *offset = d - k->tgt->tgt_start + k->tgt_offset; + void *h = k->host_start + offset; + size_t h_size = k->host_end - k->host_start; gomp_mutex_unlock (&acc_dev->lock); - - acc_unmap_data ((void *)(k->host_start + offset)); + /* PR92503 "[OpenACC] Behavior of 'acc_free' if the memory space is still + used in a mapping". */ + gomp_fatal ("refusing to free device memory space at %p that is still" + " mapped at [%p,+%d]", + d, h, (int) h_size); } else gomp_mutex_unlock (&acc_dev->lock); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c new file mode 100644 index 00000000000..4fc6068ba98 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c @@ -0,0 +1,28 @@ +/* Verify that we refuse 'acc_free', after 'acc_map_data'. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +int +main () +{ + const int N = 108; + + char *h = (char *) malloc (N); + void *d = acc_malloc (N - 10); + if (!d) + abort (); + acc_map_data (h, d, N - 19); + + fprintf (stderr, "CheCKpOInT\n"); + acc_free (d); + + return 0; +} + +/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } + { dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+89\\\]" } + { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c new file mode 100644 index 00000000000..3f6a8e57174 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c @@ -0,0 +1,27 @@ +/* Verify that we refuse 'acc_free', after 'acc_create'. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +int +main () +{ + const int N = 108; + + char *h = (char *) malloc (N); + void *d = acc_create (h, N - 1); + if (!d) + abort (); + + fprintf (stderr, "CheCKpOInT\n"); + acc_free (d); + + return 0; +} + +/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } + { dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+107\\\]" } + { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c new file mode 100644 index 00000000000..9f4504809eb --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c @@ -0,0 +1,28 @@ +/* Verify that we refuse 'acc_free', inside 'host_data', after '#pragma acc enter data create'. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +int +main () +{ + const int N = 108; + + char *h = (char *) malloc (N); +#pragma acc enter data create (h[0:N - 2]) + +#pragma acc host_data use_device (h) + { + fprintf (stderr, "CheCKpOInT\n"); + acc_free (h); + } + + return 0; +} + +/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } + { dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+106\\\]" } + { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c new file mode 100644 index 00000000000..162083051cf --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c @@ -0,0 +1,28 @@ +/* Verify that we refuse 'acc_free', after '#pragma acc enter data create'. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +int +main () +{ + const int N = 108; + + char *h = (char *) malloc (N); +#pragma acc enter data create (h[0:N - 3]) + void *d = acc_deviceptr (h); + if (!d) + abort (); + + fprintf (stderr, "CheCKpOInT\n"); + acc_free (d); + + return 0; +} + +/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } + { dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+105\\\]" } + { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c new file mode 100644 index 00000000000..bbf44319687 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c @@ -0,0 +1,31 @@ +/* Verify that we refuse 'acc_free', inside 'host_data', inside 'data'. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +int +main () +{ + const int N = 108; + + char *h = (char *) malloc (N); +#pragma acc data create (h[0:N - 44]) + { +#pragma acc host_data use_device (h) + { + fprintf (stderr, "CheCKpOInT\n"); + acc_free (h); + } + } + + return 0; +} + +/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } + TODO PR92877 + { dg-output "libgomp: cuMemGetAddressRange_v2 error: named symbol not found" { target openacc_nvidia_accel_selected } } + { dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+64\\\]" { xfail *-*-* } } + { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c new file mode 100644 index 00000000000..6212f9eae47 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c @@ -0,0 +1,32 @@ +/* Verify that we refuse 'acc_free', inside 'data'. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +int +main () +{ + const int N = 108; + + char *h = (char *) malloc (N); +#pragma acc data create (h[0:N - 21]) + { + void *d = acc_deviceptr (h); + if (!d) + abort (); + + fprintf (stderr, "CheCKpOInT\n"); + acc_free (d); + } + + return 0; +} + +/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } + TODO PR92877 + { dg-output "libgomp: cuMemGetAddressRange_v2 error: named symbol not found" { target openacc_nvidia_accel_selected } } + { dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+87\\\]" { xfail *-*-* } } + { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c index d36a2f1c304..b0a96348c3a 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c @@ -103,7 +103,10 @@ main (int argc, char **argv) if (acc_is_present (&b[0], (N * sizeof (float)))) abort (); - acc_free (d); + acc_delete (&a[0], N * sizeof (float)); + + if (acc_is_present (&a[0], N * sizeof (float))) + abort (); for (i = 0; i < N; i++) { @@ -162,7 +165,7 @@ main (int argc, char **argv) if (!acc_is_present (&b[0], (N * sizeof (float)))) abort (); - acc_free (d); + acc_delete (&b[0], N * sizeof (float)); if (acc_is_present (&b[0], (N * sizeof (float)))) abort (); @@ -557,7 +560,10 @@ main (int argc, char **argv) if (acc_is_present (&b[0], (N * sizeof (float)))) abort (); - acc_free (d); + acc_delete (&a[0], N * sizeof (float)); + + if (acc_is_present (&a[0], N * sizeof (float))) + abort (); for (i = 0; i < N; i++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-1.c index dabc7063c68..2e3b62ebbd2 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-1.c @@ -172,13 +172,13 @@ main (int argc, char **argv) exit (EXIT_FAILURE); } + acc_delete (&h_X[0], N * sizeof (float)); + acc_delete (&h_Y1[0], N * sizeof (float)); + free (h_X); free (h_Y1); free (h_Y2); - acc_free (d_X); - acc_free (d_Y); - context_check (pctx); s = cublasDestroy (h); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c index 6a52f746dcb..6bdcfe7d429 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c @@ -182,13 +182,13 @@ main (int argc, char **argv) exit (EXIT_FAILURE); } + acc_delete (&h_X[0], N * sizeof (float)); + acc_delete (&h_Y1[0], N * sizeof (float)); + free (h_X); free (h_Y1); free (h_Y2); - acc_free (d_X); - acc_free (d_Y); - context_check (pctx); s = cublasDestroy (h); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-3.c index ccd276cd98f..8f14560ea8b 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-3.c @@ -163,13 +163,13 @@ main (int argc, char **argv) exit (EXIT_FAILURE); } + acc_delete (&h_X[0], N * sizeof (float)); + acc_delete (&h_Y1[0], N * sizeof (float)); + free (h_X); free (h_Y1); free (h_Y2); - acc_free (d_X); - acc_free (d_Y); - context_check (pctx); s = cublasDestroy (h); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c index 71365e8ed32..b403a5cf5cb 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c @@ -176,13 +176,13 @@ main (int argc, char **argv) exit (EXIT_FAILURE); } + acc_delete (&h_X[0], N * sizeof (float)); + acc_delete (&h_Y1[0], N * sizeof (float)); + free (h_X); free (h_Y1); free (h_Y2); - acc_free (d_X); - acc_free (d_Y); - context_check (pctx); s = cublasDestroy (h); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c index d6655335e21..aca4c252091 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c @@ -51,7 +51,7 @@ main (int argc, char **argv) if (acc_is_present (h, 0) != 0) abort (); - acc_free (d); + acc_delete (h, N); if (acc_is_present (h, 1) != 0) abort (); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c index ee21257c9a5..de6d38b060c 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c @@ -48,7 +48,7 @@ main (int argc, char **argv) abort (); } - acc_free (d); + acc_delete (h, N); for (i = 0; i < N; i++) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c index b686cc94815..93bfb99f415 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c @@ -23,7 +23,7 @@ main (int argc, char **argv) d = acc_copyin (h, N); - acc_free (d); + acc_delete (h, N); fprintf (stderr, "CheCKpOInT\n"); acc_copyout (h, N); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-91.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-91.c index e00ef4f7206..36fff089b83 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-91.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-91.c @@ -72,6 +72,8 @@ main (int argc, char **argv) if (async > (sync * 1.5)) abort (); + acc_unmap_data (h); + acc_free (d); free (h); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c index 7ebfb8a562b..4c599cda4b3 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c @@ -112,7 +112,10 @@ main (int argc, char **argv) if (acc_is_present (&b[0], (N * sizeof (float)))) abort (); - acc_free (d); + acc_delete (&a[0], N * sizeof (float)); + + if (acc_is_present (&a[0], N * sizeof (float))) + abort (); for (i = 0; i < N; i++) { @@ -177,7 +180,7 @@ main (int argc, char **argv) if (!acc_is_present (&b[0], (N * sizeof (float)))) abort (); - acc_free (d); + acc_delete (&b[0], N * sizeof (float)); if (acc_is_present (&b[0], (N * sizeof (float)))) abort (); @@ -609,7 +612,10 @@ main (int argc, char **argv) if (acc_is_present (&b[0], (N * sizeof (float)))) abort (); - acc_free (d); + acc_delete (&a[0], N * sizeof (float)); + + if (acc_is_present (&a[0], N * sizeof (float))) + abort (); for (i = 0; i < N; i++) { -- 2.17.1