From patchwork Wed Dec 11 17:25:14 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1207797 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-515712-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="cV4HlRSw"; 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 47Y3jc4x6pz9sTr for ; Thu, 12 Dec 2019 04:25:38 +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:subject:references:date:message-id:mime-version :content-type; q=dns; s=default; b=aGZSEQ5Mn5j3EUbI9UqFKaNYofmKy MwTClfD3n0bLsX9hvLEVWzVf2iob2SmePJqNcoSkjFTKEEqThhaAjoBVOXDeJ3ha 2OsVORLs9/FXeRlJduGFoHlZPEfc/lrJRfV+Yw8z/ztD3RwvFPOVa+VcO14G3TOl vRJSNcTeeOIWeU= 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:subject:references:date:message-id:mime-version :content-type; s=default; bh=LSWWsWRq3/unmw0qNLpeWZLdzvY=; b=cV4 HlRSwJumOgXkGmdySZ7CvRQajZ3lHf1KyNjlC+V9aq8lAVminPddFsp5m4W5h2R0 s3w0o0RgbX6ggx7tuYM1rtUU0JjkgF3CaEs59wvKnDJj6gPT0d/5ZvlenxdpMdzC gTmovBRz890LHnjtjYHHdU0fibNiTInhUtwswNnU= Received: (qmail 105421 invoked by alias); 11 Dec 2019 17:25:31 -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 105411 invoked by uid 89); 11 Dec 2019 17:25:31 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-19.1 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3 autolearn=ham version=3.3.1 spammy=copyin X-HELO: esa1.mentor.iphmx.com Received: from esa1.mentor.iphmx.com (HELO esa1.mentor.iphmx.com) (68.232.129.153) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 11 Dec 2019 17:25:29 +0000 IronPort-SDR: ug/Wt/RaHpulQuN0lsee0j4mn+kXUGYZYMYV5TcP/PUHhbAP3+Y7H+lX8bHZwxWQSbx46s/vXO TTwwLSgS4ehH9tqTZbM5AekwAgz4vqaI0Vq/ZaHyrGoQSXVl319BQlik83q5HqOjCVkOaEgh/X 4cScikhmgv/tE+x+2nPCBnhp+iYgBLzlQeMVLlvTdulcJCmc7ct5YZQ/HT2CmzhCpL38h9NcuP Vlh9D9+bX9HQ+AOSC81ryXfA8zjKGztYZGE0p5H1hVPJwOyoHAszNSkW4MqB3XfGb8LLKyShnq 0Dc= Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 11 Dec 2019 09:25:27 -0800 IronPort-SDR: ulnTOxZBQA5t3Wrv9UjE8gVj1BNVZji1ivQJ/AObtQFm8lAmN0+7uKNUY/sCdC6Nn7ZhCrmlRl F1Cq3ttgNWOCmVCGtxyq3FxWLBbl30+z3chj7U8HfiaGjIAQaR4jMIS5+ck5+celddstRJZ2l3 fwLH+Z6wl9yVU7xeF4fLVTku4sfaKpytaqfOcTEjRLtrlO8cNf/U5bp+dTVa0F9hqTCHki5qjC RvcgK6pL3mrY7rq4TCD9Yqo8k2gug6I7TN9ZIcqsjLvNGwmDbdqODykcqHxZ26quPqsPoQhSox j+4= From: Thomas Schwinge To: , Julian Brown Subject: [PR92843] [OpenACC] Fix dynamic reference counting for structured 'REFCOUNT_INFINITY' References: <20191003163505.49997-2-julian@codesourcery.com> User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.1 (x86_64-pc-linux-gnu) Date: Wed, 11 Dec 2019 18:25:14 +0100 Message-ID: <875zim22rp.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Hi! See attached "[PR92843] [OpenACC] Fix dynamic reference counting for structured 'REFCOUNT_INFINITY'"; committed to trunk in r279234. Grüße Thomas From 7c8ffaf54af2c8acb77f82349aac4dd68d47ad9d Mon Sep 17 00:00:00 2001 From: tschwinge Date: Wed, 11 Dec 2019 16:49:27 +0000 Subject: [PATCH] [PR92843] [OpenACC] Fix dynamic reference counting for structured 'REFCOUNT_INFINITY' libgomp/ PR libgomp/92843 * oacc-mem.c (present_create_copy, delete_copyout): Fix dynamic reference counting for structured 'REFCOUNT_INFINITY'. Add some assertions. (goacc_insert_pointer, goacc_remove_pointer): Adjust accordingly. * testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: New file. * testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Fix OpenACC. * testsuite/libgomp.oacc-c-c++-common/lib-82.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@279234 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog | 10 + libgomp/oacc-mem.c | 42 ++-- .../libgomp.oacc-c-c++-common/clauses-1.c | 16 +- .../libgomp.oacc-c-c++-common/lib-82.c | 6 +- .../libgomp.oacc-c-c++-common/nested-1.c | 10 +- .../libgomp.oacc-c-c++-common/pr92843-1.c | 179 ++++++++++++++++++ 6 files changed, 242 insertions(+), 21 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 0a5650ed438..e5fb05aea6d 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,5 +1,15 @@ 2019-12-11 Thomas Schwinge + PR libgomp/92843 + * oacc-mem.c (present_create_copy, delete_copyout): Fix dynamic + reference counting for structured 'REFCOUNT_INFINITY'. Add some + assertions. + (goacc_insert_pointer, goacc_remove_pointer): Adjust accordingly. + * testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: New file. + * testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Fix OpenACC. + * testsuite/libgomp.oacc-c-c++-common/lib-82.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise. + * oacc-parallel.c (find_pointer, GOACC_enter_exit_data): Move... * oacc-mem.c: ... here. (gomp_acc_insert_pointer, gomp_acc_remove_pointer): Rename to diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 571e0606ac8..a809d0495a6 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -543,11 +543,11 @@ present_create_copy (unsigned f, void *h, size_t s, int async) gomp_fatal ("[%p,+%d] not mapped", (void *)h, (int)s); } + assert (n->refcount != REFCOUNT_LINK); if (n->refcount != REFCOUNT_INFINITY) - { - n->refcount++; - n->dynamic_refcount++; - } + n->refcount++; + n->dynamic_refcount++; + gomp_mutex_unlock (&acc_dev->lock); } else if (!(f & FLAG_CREATE)) @@ -573,8 +573,10 @@ present_create_copy (unsigned f, void *h, size_t s, int async) tgt = gomp_map_vars_async (acc_dev, aq, mapnum, &hostaddrs, NULL, &s, &kinds, true, GOMP_MAP_VARS_OPENACC); - /* Initialize dynamic refcount. */ - tgt->list[0].key->dynamic_refcount = 1; + n = tgt->list[0].key; + assert (n->refcount == 1); + assert (n->dynamic_refcount == 0); + n->dynamic_refcount++; d = tgt->to_free; } @@ -698,12 +700,9 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname) (void *) h, (int) s, (void *) n->host_start, (int) host_size); } - if (n->refcount == REFCOUNT_INFINITY) - { - n->refcount = 0; - n->dynamic_refcount = 0; - } - if (n->refcount < n->dynamic_refcount) + assert (n->refcount != REFCOUNT_LINK); + if (n->refcount != REFCOUNT_INFINITY + && n->refcount < n->dynamic_refcount) { gomp_mutex_unlock (&acc_dev->lock); gomp_fatal ("Dynamic reference counting assert fail\n"); @@ -711,13 +710,15 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname) if (f & FLAG_FINALIZE) { - n->refcount -= n->dynamic_refcount; + if (n->refcount != REFCOUNT_INFINITY) + n->refcount -= n->dynamic_refcount; n->dynamic_refcount = 0; } else if (n->dynamic_refcount) { + if (n->refcount != REFCOUNT_INFINITY) + n->refcount--; n->dynamic_refcount--; - n->refcount--; } if (n->refcount == 0) @@ -895,6 +896,8 @@ goacc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes, splay_tree_key n; gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, *hostaddrs, *sizes); + assert (n->refcount != REFCOUNT_INFINITY + && n->refcount != REFCOUNT_LINK); gomp_mutex_unlock (&acc_dev->lock); tgt = n->tgt; @@ -917,10 +920,11 @@ goacc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes, goacc_aq aq = get_goacc_asyncqueue (async); tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds, true, GOMP_MAP_VARS_OPENACC); + splay_tree_key n = tgt->list[0].key; + assert (n->refcount == 1); + assert (n->dynamic_refcount == 0); + n->dynamic_refcount++; gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__); - - /* Initialize dynamic refcount. */ - tgt->list[0].key->dynamic_refcount = 1; } static void @@ -950,6 +954,8 @@ goacc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async, t = n->tgt; + assert (n->refcount != REFCOUNT_INFINITY + && n->refcount != REFCOUNT_LINK); if (n->refcount < n->dynamic_refcount) { gomp_mutex_unlock (&acc_dev->lock); @@ -963,8 +969,8 @@ goacc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async, } else if (n->dynamic_refcount) { - n->dynamic_refcount--; n->refcount--; + n->dynamic_refcount--; } gomp_mutex_unlock (&acc_dev->lock); 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 b0a96348c3a..8f01d5f32f8 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c @@ -469,7 +469,9 @@ main (int argc, char **argv) if (!acc_is_present (c, (N * sizeof (float)))) abort (); - acc_copyout (b, N * sizeof (float)); + d = (float *) acc_deviceptr (b); + + acc_memcpy_from_device (b, d, N * sizeof (float)); for (i = 0; i < N; i++) { @@ -485,10 +487,22 @@ main (int argc, char **argv) if (acc_is_present (a, N * sizeof (float))) abort (); + d = (float *) acc_deviceptr (b); + + acc_unmap_data (b); + + if (acc_is_present (b, N * sizeof (float))) + abort (); + + acc_free (d); + d = (float *) acc_deviceptr (c); acc_unmap_data (c); + if (acc_is_present (c, N * sizeof (float))) + abort (); + acc_free (d); for (i = 0; i < N; i++) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-82.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-82.c index be30a7f28ac..9cf73b31964 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-82.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-82.c @@ -120,9 +120,13 @@ main (int argc, char **argv) for (i = 0; i < N; i++) { - acc_copyout (a[i], nbytes); + acc_memcpy_from_device (a[i], d_a[i], nbytes); if (*a[i] != i) abort (); + + acc_unmap_data (a[i]); + + acc_free (d_a[i]); } free (streams); 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 4c599cda4b3..3e8b42527cf 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c @@ -517,7 +517,9 @@ main (int argc, char **argv) if (!acc_is_present (c, (N * sizeof (float)))) abort (); - acc_copyout (b, N * sizeof (float)); + d = (float *) acc_deviceptr (b); + + acc_memcpy_from_device (b, d, N * sizeof (float)); for (i = 0; i < N; i++) { @@ -534,6 +536,12 @@ main (int argc, char **argv) acc_free (d); + d = (float *) acc_deviceptr (b); + + acc_unmap_data (b); + + acc_free (d); + d = (float *) acc_deviceptr (c); acc_unmap_data (c); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c new file mode 100644 index 00000000000..db5b35b08d9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c @@ -0,0 +1,179 @@ +/* Verify that 'acc_copyout' etc. is a no-op if there's still a structured + reference count. */ + +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + + +const int c0 = 58; +const int c1 = 81; + +static void +assign_array (char *array, size_t size, char value) +{ + for (size_t i = 0; i < size; ++i) + array[i] = value; +} + +static void +verify_array (const char *array, size_t size, char value) +{ + for (size_t i = 0; i < size; ++i) + assert (array[i] == value); +} + + +float global_var; +#pragma acc declare create (global_var) + +static void +test_acc_declare () +{ + assert (acc_is_present (&global_var, sizeof global_var)); + + global_var = c0; +#pragma acc update device (global_var) + + global_var = c1; + acc_copyout (&global_var, sizeof global_var); + assert (acc_is_present (&global_var, sizeof global_var)); + assert (global_var == c1); + + global_var = c1; + acc_copyout_finalize (&global_var, sizeof global_var); + assert (acc_is_present (&global_var, sizeof global_var)); + assert (global_var == c1); + + void *global_var_d_p = acc_deviceptr (&global_var); + assert (global_var_d_p); + + void *d_p = acc_copyin (&global_var, sizeof global_var); + assert (d_p == global_var_d_p); + + acc_copyout (&global_var, sizeof global_var); + assert (acc_is_present (&global_var, sizeof global_var)); + + d_p = acc_copyin (&global_var, sizeof global_var); + assert (d_p == global_var_d_p); + + d_p = acc_copyin (&global_var, sizeof global_var); + assert (d_p == global_var_d_p); + + global_var = c1; + acc_copyout_finalize (&global_var, sizeof global_var); + assert (acc_is_present (&global_var, sizeof global_var)); + assert (global_var == c1); + + global_var = c1; + acc_copyout (&global_var, sizeof global_var); + assert (acc_is_present (&global_var, sizeof global_var)); + assert (global_var == c1); +} + + +static void +test_acc_map_data () +{ + const int N = 801; + + char *h = (char *) malloc (N); + assert (h); + void *d = acc_malloc (N); + assert (d); + acc_map_data (h, d, N); + assert (acc_is_present (h, N)); + + assign_array (h, N, c0); +#pragma acc update device (h[0:N]) + + assign_array (h, N, c1); +#pragma acc exit data copyout (h[0:N]) + assert (acc_is_present (h, N)); + verify_array (h, N, c1); + + assign_array (h, N, c1); +#pragma acc exit data copyout (h[0:N]) finalize + assert (acc_is_present (h, N)); + verify_array (h, N, c1); + +#pragma acc enter data copyin (h[0:N]) + + assign_array (h, N, c1); +#pragma acc exit data copyout (h[0:N]) + assert (acc_is_present (h, N)); + verify_array (h, N, c1); + +#pragma acc enter data copyin (h[0:N]) + +#pragma acc enter data copyin (h[0:N]) + + assign_array (h, N, c1); +#pragma acc exit data copyout (h[0:N]) finalize + assert (acc_is_present (h, N)); + verify_array (h, N, c1); + + assign_array (h, N, c1); +#pragma acc exit data copyout (h[0:N]) + assert (acc_is_present (h, N)); + verify_array (h, N, c1); +} + + +static void +test_acc_data () +{ +#define N 23 + char h[N]; + + assign_array (h, N, c0); +#pragma acc data copyin (h) + { + assert (acc_is_present (h, sizeof h)); + + assign_array (h, N, c1); + acc_copyout_finalize (h, sizeof h); + assert (acc_is_present (h, sizeof h)); + verify_array (h, N, c1); + + assign_array (h, N, c1); + acc_copyout (h, sizeof h); + assert (acc_is_present (h, sizeof h)); + verify_array (h, N, c1); + + acc_copyin (h, sizeof h); + + assign_array (h, N, c1); + acc_copyout (h, sizeof h); + assert (acc_is_present (h, sizeof h)); + verify_array (h, N, c1); + + acc_copyin (h, sizeof h); + + acc_copyin (h, sizeof h); + + assign_array (h, N, c1); + acc_copyout_finalize (h, sizeof h); + assert (acc_is_present (h, sizeof h)); + verify_array (h, N, c1); + + assign_array (h, N, c1); + acc_copyout (h, sizeof h); + assert (acc_is_present (h, sizeof h)); + verify_array (h, N, c1); + } +#undef N +} + + +int +main () +{ + test_acc_declare (); + test_acc_map_data (); + test_acc_data (); + + return 0; +} -- 2.17.1