From patchwork Fri Jan 17 21:18:19 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1224977 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-517620-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.a=rsa-sha1 header.s=default header.b=qj4LM7Ei; 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 47zv8G1zzJz9sR1 for ; Sat, 18 Jan 2020 08:19:26 +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:in-reply-to:references :mime-version:content-transfer-encoding:content-type; q=dns; s= default; b=oH98Ff0POQIu9xKUVO+QVNmvTCzSxIiBSg8qibS/d3+txr2kkJ9Zg fkKlBJbDzQWt5FYd0h27wh99LzGwmG4akh5qACjpu7KEBlQedFtJrBjGZRmuCTZO AjPAnaeRQxFDQXRWz2itX3kZuBeC5EVTYcjCDR/fUYik5UzBFDbK2M= 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:in-reply-to:references :mime-version:content-transfer-encoding:content-type; s=default; bh=cdhfYNkNKKi8+xVpuiZed2G1QOw=; b=qj4LM7EiMcK9oVowtrFZ4i8Oj3D5 FvJ1rdEhqQQ6z/WtKPEvcmDFRMCAP2fWLSMR8c3Fgua7wIHnEtFfQicRu3HXpUcm nCT62qnsm5YFPFdMjwIXqtnsTNYgRO2GiEOsfuUmDDyzLwGE0jSyVG5l67hoYHib h6DSeEBBo60axOA= Received: (qmail 128416 invoked by alias); 17 Jan 2020 21:18:53 -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 128341 invoked by uid 89); 17 Jan 2020 21:18:47 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-24.3 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.1 spammy= 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; Fri, 17 Jan 2020 21:18:43 +0000 IronPort-SDR: gmJlNjesrFVdIIURErabP7sQhp8DmemIor9Y9qECytxz08cgAEMibFxCyuPqV6vp0Q2MwhZ7ST mIQFD34nOEFFc61bPFCWjsO1IZMjjjCRFfYn+8Dl8c0/vjsNasCloxp0vPzAj6xkSuT4k2MN/X /L3j7yA3gW9PBO8QINRu7jtMV8lONjkrOiONzcsEk2KKhF8jCdVA+fKs3Win8J/M/uNuvHiUmr iXA7gKcsaSXIcOURIrXRsEjS72Bb+QzVtREICYhIIcWlrQOQB9mNh9KiZrOkZhQGpzHU17O3nb 6hs= Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 17 Jan 2020 13:18:43 -0800 IronPort-SDR: EI7CuJ7gVnCqXEOpzQxN5r2FGE8BpTtTEbUtu28RdYmgLsrxthytQRVQimI/PrxNctWlMUFZjL 9XjtAmFsK/DA== From: Julian Brown To: CC: Thomas Schwinge , Jakub Jelinek Subject: [PATCH 1/3] Introduce dynamic data mapping sentinel for OpenACC Date: Fri, 17 Jan 2020 13:18:19 -0800 Message-ID: <666bd70af1514e8a3d80fa9ad1ef114bbc2b3a80.1579292772.git.julian@codesourcery.com> In-Reply-To: References: MIME-Version: 1.0 X-IsSubscribed: yes This patch provides a way to distinguish target_mem_descs that arise from "enter data" operations from those that arise from structured OpenACC data blocks. In that way, we can implement the equivalent of the "no-op" behaviour of decrementing a dynamic reference count that is already zero for some given variable, as described in the OpenACC 2.6 spec. We do this by re-using the "prev" field of the target_mem_desc (currently unused for dynamic data mappings) to store a special sentinel value. Several new tests are added, both for cases that now work, and for diagnostics for cases that do not. Tested alongside other patches in this series with offloading to NVPTX. OK? Thanks, Julian ChangeLog PR libgomp/92843 libgomp/ * libgomp.h (target_mem_desc): Update comment for prev field. * oacc-int.h (goacc_mark_dynamic): Add prototype. * oacc-mem.c (dyn_tgt_sentinel): New static global. (goacc_mark_dynamic): New function. (goacc_enter_datum, goacc_enter_data_internal): Call goacc_mark_dynamic on non-NULL target_mem_desc return from gomp_map_vars_async. (goacc_exit_datum, goacc_exit_data_internal): Check target_mem_desc for sentinel value on structural refcount decrement. * target.c (gomp_unmap_vars_internal): Re-use target_mem_desc for "structural" data mapping for extending dynamic mapping beyond the end of a structured block when possible. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c: Likewise. --- libgomp/libgomp.h | 3 +- libgomp/oacc-int.h | 2 + libgomp/oacc-mem.c | 53 ++++- libgomp/target.c | 56 +++++- .../static-dynamic-lifetimes-2-lib.c | 3 + .../static-dynamic-lifetimes-2.c | 166 ++++++++++++++++ .../static-dynamic-lifetimes-3-lib.c | 3 + .../static-dynamic-lifetimes-3.c | 183 ++++++++++++++++++ .../static-dynamic-lifetimes-4-lib.c | 6 + .../static-dynamic-lifetimes-4.c | 71 +++++++ .../static-dynamic-lifetimes-5-lib.c | 6 + .../static-dynamic-lifetimes-5.c | 63 ++++++ 12 files changed, 596 insertions(+), 19 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 15a1394c16d..bbab4f9f34f 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -975,7 +975,8 @@ struct target_mem_desc { uintptr_t tgt_end; /* Handle to free. */ void *to_free; - /* Previous target_mem_desc. */ + /* Previous target_mem_desc. Also used in OpenACC to indicate that this + target_mem_desc is used only for an "enter data" mapping. */ struct target_mem_desc *prev; /* Number of items in following list. */ size_t list_count; diff --git a/libgomp/oacc-int.h b/libgomp/oacc-int.h index 3c2c9b84b2f..bb67188c3e9 100644 --- a/libgomp/oacc-int.h +++ b/libgomp/oacc-int.h @@ -165,6 +165,8 @@ bool _goacc_profiling_setup_p (struct goacc_thread *, void goacc_profiling_dispatch (acc_prof_info *, acc_event_info *, acc_api_info *); +extern void goacc_mark_dynamic (struct target_mem_desc *); + #ifdef HAVE_ATTRIBUTE_VISIBILITY # pragma GCC visibility pop #endif diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index bd1a99d9277..45ab2b169d7 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -550,6 +550,24 @@ acc_unmap_data (void *h) } } +/* Indicate (via storing its address in the "prev" field) a target_mem_desc + that is used for an "enter data" mapping. */ +const static struct target_mem_desc dyn_tgt_sentinel; + +/* Mark TGT as the "initial" target_mem_desc created by a dynamic data mapping + (acc_create, acc_copyin or an "enter data" directive). For such mappings, + to start with, we have a splay tree key with a reference count of 1 and a + virtual reference count of 0 (linking to this target_mem_desc). Without + this marking, such a mapping is indistinguishable from a target_mem_desc + created by e.g. a lexically-scoped "acc data" region, but the difference is + important if acc_copyout, acc_delete (etc.) or an "exit data" directive is + used to end the data lifetime. */ + +void +goacc_mark_dynamic (struct target_mem_desc *tgt) +{ + tgt->prev = (struct target_mem_desc *) &dyn_tgt_sentinel; +} /* Enter dynamic mapping for a single datum. Return the device pointer. */ @@ -613,8 +631,14 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async) goacc_aq aq = get_goacc_asyncqueue (async); - gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds, - true, GOMP_MAP_VARS_OPENACC_ENTER_DATA); + struct target_mem_desc *tgt; + tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, + kinds, true, GOMP_MAP_VARS_OPENACC_ENTER_DATA); + + /* Mark non-NULL target_mem_descs returned here specially: see comment in + goacc_exit_datum. */ + if (tgt) + goacc_mark_dynamic (tgt); gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, hostaddrs[0], sizes[0]); @@ -756,7 +780,15 @@ goacc_exit_datum (void *h, size_t s, unsigned short kind, int async) n->refcount--; n->virtual_refcount--; } - else if (n->refcount > 0 && n->refcount != REFCOUNT_INFINITY) + /* An initial "enter data" mapping might create a target_mem_desc (in + gomp_map_vars_async via goacc_enter_datum). In that case we have a + structural reference count but a zero virtual reference count: we + nevertheless want to do the "exit data" operation here. Detect the + special case using a sentinel value stored in the "prev" field, which is + otherwise unused for dynamic data mappings. */ + else if (n->refcount > 0 + && n->refcount != REFCOUNT_INFINITY + && n->tgt->prev == &dyn_tgt_sentinel) n->refcount--; if (n->refcount == 0) @@ -1081,11 +1113,12 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, dump_mappings ((group_last - i) + 1, &hostaddrs[i], &sizes[i], &kinds[i]); #endif - gomp_map_vars_async (acc_dev, aq, - (group_last - i) + 1, - &hostaddrs[i], NULL, - &sizes[i], &kinds[i], true, - GOMP_MAP_VARS_OPENACC_ENTER_DATA); + struct target_mem_desc *tgt; + tgt = gomp_map_vars_async (acc_dev, aq, (group_last - i) + 1, + &hostaddrs[i], NULL, &sizes[i], &kinds[i], + true, GOMP_MAP_VARS_OPENACC_ENTER_DATA); + if (tgt) + goacc_mark_dynamic (tgt); i = group_last; } @@ -1196,7 +1229,9 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, n->refcount--; n->virtual_refcount--; } - else if (n->refcount > 0 && n->refcount != REFCOUNT_INFINITY) + else if (n->refcount > 0 + && n->refcount != REFCOUNT_INFINITY + && n->tgt->prev == &dyn_tgt_sentinel) n->refcount--; if (copyfrom diff --git a/libgomp/target.c b/libgomp/target.c index 825213f40ec..fb423ced144 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1629,6 +1629,8 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom, k->refcount == 1, NULL); } + bool have_virtual_refs = false, all_refs_virtual = true; + for (i = 0; i < tgt->list_count; i++) { splay_tree_key k = tgt->list[i].key; @@ -1636,21 +1638,21 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom, continue; bool do_unmap = false; - if (k->tgt == tgt - && k->virtual_refcount > 0 - && k->refcount != REFCOUNT_INFINITY) - { - k->virtual_refcount--; - k->refcount--; - } - else if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY) + if (k->refcount > 1 && k->refcount != REFCOUNT_INFINITY) k->refcount--; else if (k->refcount == 1) { k->refcount--; - do_unmap = true; + if (k->virtual_refcount == 0) + do_unmap = true; } + if (k->virtual_refcount > 0 && k->refcount == k->virtual_refcount) + have_virtual_refs = true; + + if (k->refcount != k->virtual_refcount) + all_refs_virtual = false; + if ((do_unmap && do_copyfrom && tgt->list[i].copy_from) || tgt->list[i].always_copy_from) gomp_copy_dev2host (devicep, aq, @@ -1670,6 +1672,42 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom, } } + if (have_virtual_refs) + { + /* If we have a construct such as this: + + #pragma acc data copy(var1) + { + #pragma acc enter data copyin(var1) + } + + The dynamic data lifetime entered in the middle of the static + data lifetime extends beyond the static lifetime. Adjust + references and the target descriptor here (the end of the static + region) to make it seem like we did "enter data" on the data to + start with. + + We can't do this adjustment if the data construct refers to other + variables too. */ + if (!all_refs_virtual) + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("cannot handle create/copyin/'enter data' within data " + "region"); + } + + for (i = 0; i < tgt->list_count; i++) + { + splay_tree_key k = tgt->list[i].key; + if (k == NULL || k->virtual_refcount == 0) + continue; + + if (k->refcount == k->virtual_refcount) + k->virtual_refcount--; + } + goacc_mark_dynamic (tgt); + } + if (aq) devicep->openacc.async.queue_callback_func (aq, gomp_unref_tgt_void, (void *) tgt); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c new file mode 100644 index 00000000000..84f41a49dfd --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c @@ -0,0 +1,3 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "static-dynamic-lifetimes-2.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c new file mode 100644 index 00000000000..d3c6f5192d8 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c @@ -0,0 +1,166 @@ +/* Test nested dynamic/static data mappings. */ + +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#define SIZE 1024 + +void +f1 (void) +{ + char *block1 = (char *) malloc (SIZE); + +#pragma acc data copy(block1[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyin (block1, SIZE); + acc_copyout (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#pragma acc exit data copyout(block1[0:SIZE]) +#endif + } + + assert (!acc_is_present (block1, SIZE)); + + free (block1); +} + +void +f2 (void) +{ + char *block1 = (char *) malloc (SIZE); + +#ifdef OPENACC_API + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE]) + { + } + +#ifdef OPENACC_API + acc_copyout (block1, SIZE); +#else +#pragma acc exit data copyout(block1[0:SIZE]) +#endif + + assert (!acc_is_present (block1, SIZE)); + + free (block1); +} + +void +f3 (void) +{ + char *block1 = (char *) malloc (SIZE); + +#pragma acc data copy(block1[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyin (block1, SIZE); + acc_copyin (block1, SIZE); + acc_copyout (block1, SIZE); + acc_copyout (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#pragma acc enter data copyin(block1[0:SIZE]) +#pragma acc exit data copyout(block1[0:SIZE]) +#pragma acc exit data copyout(block1[0:SIZE]) +#endif + } + + assert (!acc_is_present (block1, SIZE)); + + free (block1); +} + +void +f4 (void) +{ + char *block1 = (char *) malloc (SIZE); + +#pragma acc data copy(block1[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyin (block1, SIZE); + acc_copyout (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#pragma acc exit data copyout(block1[0:SIZE]) +#endif + } + +#ifdef OPENACC_API + acc_copyout (block1, SIZE); +#else +#pragma acc exit data copyout(block1[0:SIZE]) +#endif + } + + assert (!acc_is_present (block1, SIZE)); + + free (block1); +} + +void +f5 (void) +{ + char *block1 = (char *) malloc (SIZE); + +#ifdef OPENACC_API + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#endif +#pragma acc data copy(block1[0:SIZE]) + { + } +#ifdef OPENACC_API + acc_copyout (block1, SIZE); +#else +#pragma acc exit data copyout(block1[0:SIZE]) +#endif + } +#ifdef OPENACC_API + acc_copyout (block1, SIZE); +#else +#pragma acc exit data copyout(block1[0:SIZE]) +#endif + + assert (!acc_is_present (block1, SIZE)); + + free (block1); +} + +int +main (int argc, char *argv[]) +{ + f1 (); + f2 (); + f3 (); + f4 (); + f5 (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c new file mode 100644 index 00000000000..d9e76c600f0 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c @@ -0,0 +1,3 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "static-dynamic-lifetimes-3.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c new file mode 100644 index 00000000000..59501864398 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c @@ -0,0 +1,183 @@ +/* Test nested dynamic/static data mappings (multiple blocks on data + regions). */ + +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#define SIZE 1024 + +void +f1 (void) +{ + char *block1 = (char *) malloc (SIZE); + char *block2 = (char *) malloc (SIZE); + +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyin (block1, SIZE); + acc_copyout (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#pragma acc exit data copyout(block1[0:SIZE]) +#endif + } + + assert (!acc_is_present (block1, SIZE)); + assert (!acc_is_present (block2, SIZE)); + + free (block1); + free (block2); +} + +void +f2 (void) +{ + char *block1 = (char *) malloc (SIZE); + char *block2 = (char *) malloc (SIZE); + +#ifdef OPENACC_API + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE]) + { + } + +#ifdef OPENACC_API + acc_copyout (block1, SIZE); +#else +#pragma acc exit data copyout(block1[0:SIZE]) +#endif + + assert (!acc_is_present (block1, SIZE)); + assert (!acc_is_present (block2, SIZE)); + + free (block1); + free (block2); +} + +void +f3 (void) +{ + char *block1 = (char *) malloc (SIZE); + char *block2 = (char *) malloc (SIZE); + +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyin (block1, SIZE); + acc_copyin (block2, SIZE); + acc_copyout (block2, SIZE); + acc_copyout (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#pragma acc enter data copyin(block2[0:SIZE]) +#pragma acc exit data copyout(block2[0:SIZE]) +#pragma acc exit data copyout(block1[0:SIZE]) +#endif + } + + assert (!acc_is_present (block1, SIZE)); + assert (!acc_is_present (block2, SIZE)); + + free (block1); + free (block2); +} + +void +f4 (void) +{ + char *block1 = (char *) malloc (SIZE); + char *block2 = (char *) malloc (SIZE); + +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyin (block2, SIZE); + acc_copyout (block2, SIZE); +#else +#pragma acc enter data copyin(block2[0:SIZE]) +#pragma acc exit data copyout(block2[0:SIZE]) +#endif + } +#ifdef OPENACC_API + acc_copyout (block1, SIZE); +#else +#pragma acc exit data copyout(block1[0:SIZE]) +#endif + } + + assert (!acc_is_present (block1, SIZE)); + assert (!acc_is_present (block2, SIZE)); + + free (block1); + free (block2); +} + +void +f5 (void) +{ + char *block1 = (char *) malloc (SIZE); + char *block2 = (char *) malloc (SIZE); + +#ifdef OPENACC_API + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyin (block2, SIZE); +#else +#pragma acc enter data copyin(block2[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE]) + { + } +#ifdef OPENACC_API + acc_copyout (block2, SIZE); +#else +#pragma acc exit data copyout(block2[0:SIZE]) +#endif + } + +#ifdef OPENACC_API + acc_copyout (block1, SIZE); +#else +#pragma acc exit data copyout(block1[0:SIZE]) +#endif + + assert (!acc_is_present (block1, SIZE)); + assert (!acc_is_present (block2, SIZE)); + + free (block1); + free (block2); +} + +int +main (int argc, char *argv[]) +{ + f1 (); + f2 (); + f3 (); + f4 (); + f5 (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c new file mode 100644 index 00000000000..77bcd9e8dd8 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c @@ -0,0 +1,6 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "static-dynamic-lifetimes-4.c" + +/* { dg-output "libgomp: cannot handle create/copyin/.enter data. within data region" } */ +/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c new file mode 100644 index 00000000000..0d9f52febdb --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c @@ -0,0 +1,71 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#define SIZE 1024 + +int +main (int argc, char *argv[]) +{ + char *block1 = (char *) malloc (SIZE); + char *block2 = (char *) malloc (SIZE); + char *block3 = (char *) malloc (SIZE); + + /* Doing this twice ensures that we have a non-zero virtual refcount. Make + sure that works too. */ +#ifdef OPENACC_API + acc_copyin (block1, SIZE); + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE], block3[0:SIZE]) + { + /* The first copyin of block2 is the enclosing data region. This + "enter data" should make it live beyond the end of this region. */ +#ifdef OPENACC_API + acc_copyin (block2, SIZE); + /* Error output checked in static-dynamic-lifetimes-4-lib.c. */ +#else +#pragma acc enter data copyin(block2[0:SIZE]) + /* ...except that doesn't work at present because it would mean the dynamic + data region would get entangled with the static data region's + target_mem_desc that has mappings for each of block1, block2 and block3. + Check for runtime error. */ + /* { dg-output "libgomp: cannot handle create/copyin/.enter data. within data region" } */ + /* { dg-shouldfail "" } */ +#endif + } + + assert (acc_is_present (block1, SIZE)); + assert (acc_is_present (block2, SIZE)); + assert (!acc_is_present (block3, SIZE)); + +#ifdef OPENACC_API + acc_copyout (block1, SIZE); + assert (acc_is_present (block1, SIZE)); + acc_copyout (block1, SIZE); + assert (!acc_is_present (block1, SIZE)); + + acc_copyout (block2, SIZE); + assert (!acc_is_present (block2, SIZE)); +#else +#pragma acc exit data copyout(block1[0:SIZE]) + assert (acc_is_present (block1, SIZE)); +#pragma acc exit data copyout(block1[0:SIZE]) + assert (!acc_is_present (block1, SIZE)); + +#pragma acc exit data copyout(block2[0:SIZE]) + assert (!acc_is_present (block2, SIZE)); +#endif + + free (block1); + free (block2); + free (block3); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c new file mode 100644 index 00000000000..dcf4da6b660 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c @@ -0,0 +1,6 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "static-dynamic-lifetimes-5.c" + +/* { dg-output "libgomp: cannot handle create/copyin/.enter data. within data region" } */ +/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c new file mode 100644 index 00000000000..062ca74f2ab --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c @@ -0,0 +1,63 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#define SIZE 1024 + +int +main (int argc, char *argv[]) +{ + char *block1 = (char *) malloc (SIZE); + char *block2 = (char *) malloc (SIZE); + char *block3 = (char *) malloc (SIZE); + +#ifdef OPENACC_API + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE], block3[0:SIZE]) + { + /* The first copyin of block2 is the enclosing data region. This + "enter data" should make it live beyond the end of this region. */ +#ifdef OPENACC_API + acc_copyin (block2, SIZE); + /* Error output checked in static-dynamic-lifetimes-5-lib.c. */ +#else +#pragma acc enter data copyin(block2[0:SIZE]) + /* ...except that doesn't work at present because it would mean the dynamic + data region would get entangled with the static data region's + target_mem_desc that has mappings for each of block1, block2 and block3. + Check for runtime error. */ + /* { dg-output "libgomp: cannot handle create/copyin/.enter data. within data region" } */ + /* { dg-shouldfail "" } */ +#endif + } + + assert (acc_is_present (block1, SIZE)); + assert (acc_is_present (block2, SIZE)); + assert (!acc_is_present (block3, SIZE)); + +#ifdef OPENACC_API + acc_copyout (block1, SIZE); + assert (!acc_is_present (block1, SIZE)); + + acc_copyout (block2, SIZE); + assert (!acc_is_present (block2, SIZE)); +#else +#pragma acc exit data copyout(block1[0:SIZE]) + assert (!acc_is_present (block1, SIZE)); + +#pragma acc exit data copyout(block2[0:SIZE]) + assert (!acc_is_present (block2, SIZE)); +#endif + + free (block1); + free (block2); + free (block3); + + return 0; +} From patchwork Fri Jan 17 21:18:20 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1224976 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-517619-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.a=rsa-sha1 header.s=default header.b=uPmgVdvG; 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 47zv805HdMz9sR1 for ; Sat, 18 Jan 2020 08:19:12 +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:in-reply-to:references :mime-version:content-transfer-encoding:content-type; q=dns; s= default; b=QefERj62EYpO2oH7WBzjdK61nc/H5P8jee8M+vxC0W2/uvPTnMw1D SbyeKC6GyrOzG2Wlfnry0q/KoyQgveZxsFQB22EWYO5MzeeVRX0w/htPF+Ev6kL6 mu3mRBB6NbwvKEhBsq5Q6xRu3qMNVppyhOUSX3RgIMAq3jGT/EAx3A= 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:in-reply-to:references :mime-version:content-transfer-encoding:content-type; s=default; bh=rBpRqqGPmWZFEX+t1xAjYuN4A9I=; b=uPmgVdvGsdEvt5rrdxd6kau4U1Lm sXW/BYkjk0Q7zj21Y+Kly8namOlRvGcFrliGKMTSTWjhLclNtwNeYymZt/5+2FWe rOC5sq1F9y9ZSI5UOYAUB34JrNVGxqBOF0Z8Juh0qdC1v1KZcAi7HX53QS/m7HC4 z4g+52W4FvF1asM= Received: (qmail 128402 invoked by alias); 17 Jan 2020 21:18:53 -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 128343 invoked by uid 89); 17 Jan 2020 21:18:48 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-24.3 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= 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; Fri, 17 Jan 2020 21:18:46 +0000 IronPort-SDR: IPiH4cQ0gm3JkdOjf4vq1qxlY/MZvc4r8EfCxnAi+z0RkUYK7E3LE2ga3uVW2StfD+e8cUbOZv JGo+UyOTJICemvE6GRFO9JXq0lrNeGxD0qNb37P2j/oxh4uuFZ0R+afc3ZfdWnwMJmKJQpmbID yJzlHFpgFCD4rLKxRdb0AyC2TVOjeP7bdjVqtgqv3TVsbNwRan2d4bOnaBIbncdINChZEQJQCH OoKiRDYUtLI2fDfCNybGHE1/6rh6zxiYzki/UpO09pq1nlrLgpBdn130dndKsqaTunA/pq8uaz JVU= Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 17 Jan 2020 13:18:45 -0800 IronPort-SDR: swdSgnJLouro1gZhixSoV9wFZVgF/jcZ4CiCN1rXHNu/1dL1OoQ2dN0aLewou9+HE8sEdKlXGO v3U/QJ+swZNA== From: Julian Brown To: CC: Thomas Schwinge , Jakub Jelinek Subject: [PATCH 2/3] Don't copy back vars mapped with acc_map_data Date: Fri, 17 Jan 2020 13:18:20 -0800 Message-ID: <4bbb4970cd48424873573a50e627786cc9cf3378.1579292772.git.julian@codesourcery.com> In-Reply-To: References: MIME-Version: 1.0 X-IsSubscribed: yes This patch prevents "exit data" directives from copying back data that was mapped with an acc_map_data API call. This matches the behaviour expected by the pr92843-1.c test, and together with the previous patch in this series, allows that test to pass (with no other regressions). Tested alongside other patches in this series with offloading to NVPTX (with and without the third & final patch). OK? Thanks, Julian ChangeLog PR libgomp/92843 libgomp/ * oacc-mem.c (goacc_exit_data_internal): Don't copy-back data mapped with acc_map_data on an "exit data" directive. * testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Remove XFAIL. Add explanatory comment. --- libgomp/oacc-mem.c | 1 + libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c | 4 +++- 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 45ab2b169d7..783e7f363fb 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1235,6 +1235,7 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, n->refcount--; 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 diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c index f16c46a37bf..786a12a8504 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c @@ -1,7 +1,6 @@ /* Verify that 'acc_copyout' etc. is a no-op if there's still a structured reference count. */ -/* { dg-xfail-run-if "TODO PR92843" { *-*-* } } */ /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ #include @@ -96,6 +95,9 @@ test_acc_map_data () verify_array (h, N, c1); assign_array (h, N, c1); + /* Note that we're not expecting this (nor the copyouts below) to perform + an actual "finalize" or copyout since the data was mapped with + acc_map_data. */ #pragma acc exit data copyout (h[0:N]) finalize assert (acc_is_present (h, N)); verify_array (h, N, c1); From patchwork Fri Jan 17 21:18:21 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1224978 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-517621-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.a=rsa-sha1 header.s=default header.b=bgsW/aDZ; 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 47zv8X5PDMz9sR4 for ; Sat, 18 Jan 2020 08:19:40 +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:in-reply-to:references :mime-version:content-transfer-encoding:content-type; q=dns; s= default; b=H+E8RbJtV5XhbeGh3rzBjCx+VYPcmg5CvoeEj4pvPgKyX6zY4iQ0E HEttB5uXSb0XvGUDCb7VBTYCYSFlK/exXATOXDamOTOrikeXLd4noMae/YP/PF9d o82POGJb7JTbwyEy07K/r6IewlucnMqKX4kopbkgJEOfUuFRDnJACw= 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:in-reply-to:references :mime-version:content-transfer-encoding:content-type; s=default; bh=Y0IPkGZK3e2lwlXu+V2FsMgvwfs=; b=bgsW/aDZDWDoZiUmZMDoSey4Nhop c6sC+ODQNhpRBsa3wkRUeXv1APKr1PoRf/Fx1a+ODAjZEFjUOYOcoVQrSBCchHYu 3cWDauf2L+xCfqReIYYNSmGeZ5Sbem+UyT4N1QihwFucQmwjurAecaxF+a1mW9w7 sLCcZbAKyJm+Ifs= Received: (qmail 128567 invoked by alias); 17 Jan 2020 21:18:54 -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 128425 invoked by uid 89); 17 Jan 2020 21:18:53 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-24.3 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.1 spammy= 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; Fri, 17 Jan 2020 21:18:50 +0000 IronPort-SDR: /x5WTuKpSg+ADo9B9uBqtsT9nE6B0NAgum0t9tnWYrs6qpvDG+MbeamhXw/tGjbuypeLphnJpp i+xCSGXojMVI1e/vyhjj9DqHU3Z2bVeKtWgbFrXL/ugC7I14YLisbwjpX3Rg7ots/fkUMhQGih hsX6KKZyNTb7S0LJcn7NP+L/kj8T3QBjoCsZ1wZf/BZB0XDig/dqhOf0we7+7e17/qFibwAg0D WVqYFAKA6e8ox23UV1lSKnpM3pN5gnwdkB4KHULnVWD68xmU8TW7EhxciaEdz1CKzsYFtrHnVz SUM= Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 17 Jan 2020 13:18:49 -0800 IronPort-SDR: +x9PcVKobWR0lzjfxes0LcLUBo1MOqGG/gy/t2dO3qc/o3A/Bd/Y/6p/bynW4pD5+V0DSaBwHD 0KI39zWB0Y3g== From: Julian Brown To: CC: Thomas Schwinge , Jakub Jelinek Subject: [PATCH 3/3] OpenACC dynamic data lifetimes ending within structured blocks Date: Fri, 17 Jan 2020 13:18:21 -0800 Message-ID: <4673a5070087e465f6dd123715d409b35b875ca1.1579292772.git.julian@codesourcery.com> In-Reply-To: References: MIME-Version: 1.0 X-IsSubscribed: yes This patch adds a new function to logically decrement the "dynamic reference counter" for a mapped OpenACC variable, and handles some cases in which that counter drops to zero inside a structured data block. Previously, it's likely that at least in some cases, ending a dynamic data lifetime in this way could behave unpredictably. Several new test cases are included. This patch is strongly related to the previous two, but is somewhat of a separate change, and those two patches can stand alone if this one gets deferred. Tested alongside the previous patches in the series with offloading to NVPTX. OK? Thanks, Julian ChangeLog libgomp/ * oacc-mem.c (decr_dynamic_refcount): New function. (goacc_exit_datum): Call above function. (goacc_exit_data_internal): Call above function. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c: Likewise. --- libgomp/oacc-mem.c | 128 ++++++++++---- .../static-dynamic-lifetimes-1-lib.c | 3 + .../static-dynamic-lifetimes-1.c | 160 ++++++++++++++++++ .../static-dynamic-lifetimes-6-lib.c | 5 + .../static-dynamic-lifetimes-6.c | 46 +++++ .../static-dynamic-lifetimes-7-lib.c | 5 + .../static-dynamic-lifetimes-7.c | 45 +++++ .../static-dynamic-lifetimes-8-lib.c | 5 + .../static-dynamic-lifetimes-8.c | 50 ++++++ 9 files changed, 412 insertions(+), 35 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 783e7f363fb..f34ffa67079 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -725,6 +725,92 @@ acc_pcopyin (void *h, size_t s) #endif +/* Perform actions necessary to decrement the dynamic reference counter for + splay tree key N. Returns TRUE on success, or FALSE on failure (e.g. if we + hit a case we can't presently handle inside a data region). */ + +static bool +decr_dynamic_refcount (splay_tree_key n, bool finalize) +{ + if (finalize) + { + if (n->refcount != REFCOUNT_INFINITY) + n->refcount -= n->virtual_refcount; + n->virtual_refcount = 0; + } + + if (n->virtual_refcount > 0) + { + if (n->refcount != REFCOUNT_INFINITY) + n->refcount--; + n->virtual_refcount--; + } + /* An initial "enter data" mapping might create a target_mem_desc (in + gomp_map_vars_async via goacc_enter_datum or + goacc_enter_data_internal). In that case we have a structural + reference count but a zero virtual reference count: we nevertheless + want to do the "exit data" operation here. Detect the special case + using a sentinel value stored in the "prev" field, which is otherwise + unused for dynamic data mappings. */ + else if (n->refcount > 0 + && n->refcount != REFCOUNT_INFINITY + && n->tgt->prev == &dyn_tgt_sentinel) + { + n->refcount--; + /* We know n->virtual_refcount is zero here, so if we still have a + non-zero n->refcount we are ending a dynamically-scoped variable + lifetime in the middle of a static lifetime for the same variable. + If we're not careful this results in a dangling reference. Attempt + to handle this here, if only in simple cases. E.g.: + + #pragma acc enter data copyin(var) + #pragma acc data copy(var{, ...}) + { + #pragma acc exit data copyout(var) + } + + Here (the "exit data"), we reattach the relevant fields of the + previously dynamically-scoped target_mem_desc to the static data + region's target_mem_desc, hence merging the former into the latter. + The old dynamic target_mem_desc can then be freed. + + We can't deal with static data regions that refer to existing dynamic + data mappings or that introduce new static lifetimes of their own. */ + if (n->refcount > 0 + && n->tgt->list_count == 1 + && n->tgt->refcount == 1) + { + struct goacc_thread *thr = goacc_thread (); + struct target_mem_desc *tgt, *static_tgt = NULL; + for (tgt = thr->mapped_data; + tgt != NULL && static_tgt == NULL; + tgt = tgt->prev) + for (int j = 0; j < tgt->list_count; j++) + if (tgt->list[j].key == n) + { + static_tgt = tgt; + break; + } + if (!static_tgt + || static_tgt->to_free != NULL + || static_tgt->array != NULL) + return false; + static_tgt->to_free = n->tgt->to_free; + static_tgt->array = n->tgt->array; + static_tgt->tgt_start = n->tgt->tgt_start; + static_tgt->tgt_end = n->tgt->tgt_end; + static_tgt->to_free = n->tgt->to_free; + static_tgt->refcount++; + free (n->tgt); + n->tgt = static_tgt; + } + else if (n->refcount > 0) + return false; + } + + return true; +} + /* Exit a dynamic mapping for a single variable. */ static void @@ -767,29 +853,12 @@ goacc_exit_datum (void *h, size_t s, unsigned short kind, int async) bool finalize = (kind == GOMP_MAP_DELETE || kind == GOMP_MAP_FORCE_FROM); - if (finalize) - { - if (n->refcount != REFCOUNT_INFINITY) - n->refcount -= n->virtual_refcount; - n->virtual_refcount = 0; - } - if (n->virtual_refcount > 0) + if (!decr_dynamic_refcount (n, finalize)) { - if (n->refcount != REFCOUNT_INFINITY) - n->refcount--; - n->virtual_refcount--; + gomp_mutex_unlock (&acc_dev->lock); + gomp_fatal ("cannot handle delete/copyout within data region"); } - /* An initial "enter data" mapping might create a target_mem_desc (in - gomp_map_vars_async via goacc_enter_datum). In that case we have a - structural reference count but a zero virtual reference count: we - nevertheless want to do the "exit data" operation here. Detect the - special case using a sentinel value stored in the "prev" field, which is - otherwise unused for dynamic data mappings. */ - else if (n->refcount > 0 - && n->refcount != REFCOUNT_INFINITY - && n->tgt->prev == &dyn_tgt_sentinel) - n->refcount--; if (n->refcount == 0) { @@ -1216,23 +1285,12 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, if (n == NULL) continue; - if (finalize) - { - if (n->refcount != REFCOUNT_INFINITY) - n->refcount -= n->virtual_refcount; - n->virtual_refcount = 0; - } - - if (n->virtual_refcount > 0) + if (!decr_dynamic_refcount (n, finalize)) { - if (n->refcount != REFCOUNT_INFINITY) - n->refcount--; - n->virtual_refcount--; + /* The user is trying to do something too tricky for us. */ + gomp_mutex_unlock (&acc_dev->lock); + gomp_fatal ("cannot handle 'exit data' within data region"); } - else if (n->refcount > 0 - && n->refcount != REFCOUNT_INFINITY - && n->tgt->prev == &dyn_tgt_sentinel) - n->refcount--; if (copyfrom && n->refcount != REFCOUNT_INFINITY diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c new file mode 100644 index 00000000000..23c20d4fab7 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c @@ -0,0 +1,3 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "static-dynamic-lifetimes-1.c" diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c new file mode 100644 index 00000000000..a743660f53e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c @@ -0,0 +1,160 @@ +/* Test transitioning of data lifetimes between static and dynamic. */ + +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#define SIZE 1024 + +void +f1 (void) +{ + char *block1 = (char *) malloc (SIZE); + +#ifdef OPENACC_API + acc_copyin (block1, SIZE); + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + } + + assert (acc_is_present (block1, SIZE)); + +#ifdef OPENACC_API + acc_copyout (block1, SIZE); + assert (acc_is_present (block1, SIZE)); + acc_copyout (block1, SIZE); + assert (acc_is_present (block1, SIZE)); + acc_copyout (block1, SIZE); + assert (!acc_is_present (block1, SIZE)); +#else +#pragma acc exit data copyout(block1[0:SIZE]) + assert (acc_is_present (block1, SIZE)); +#pragma acc exit data copyout(block1[0:SIZE]) + assert (acc_is_present (block1, SIZE)); +#pragma acc exit data copyout(block1[0:SIZE]) + assert (!acc_is_present (block1, SIZE)); +#endif + + free (block1); +} + +void +f2 (void) +{ + char *block1 = (char *) malloc (SIZE); + +#ifdef OPENACC_API + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyout (block1, SIZE); +#else +#pragma acc exit data copyout(block1[0:SIZE]) +#endif + /* This should stay present until the end of the static data lifetime. */ + assert (acc_is_present (block1, SIZE)); + } + + assert (!acc_is_present (block1, SIZE)); + + free (block1); +} + +void +f3 (void) +{ + char *block1 = (char *) malloc (SIZE); + +#ifdef OPENACC_API + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyout (block1, SIZE); + acc_copyin (block1, SIZE); +#else +#pragma acc exit data copyout(block1[0:SIZE]) +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + assert (acc_is_present (block1, SIZE)); + } + + assert (acc_is_present (block1, SIZE)); +#ifdef OPENACC_API + acc_copyout (block1, SIZE); +#else +#pragma acc exit data copyout(block1[0:SIZE]) +#endif + assert (!acc_is_present (block1, SIZE)); + + free (block1); +} + +void +f4 (void) +{ + char *block1 = (char *) malloc (SIZE); + char *block2 = (char *) malloc (SIZE); + char *block3 = (char *) malloc (SIZE); + +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE], block3[0:SIZE]) + { + /* The first copyin of block2 is the enclosing data region. This + "enter data" should make it live beyond the end of this region. + This works, though the on-target copies of block1, block2 and block3 + will stay allocated until block2 is unmapped because they are bound + together in a single target_mem_desc. */ +#ifdef OPENACC_API + acc_copyin (block2, SIZE); +#else +#pragma acc enter data copyin(block2[0:SIZE]) +#endif + } + + assert (!acc_is_present (block1, SIZE)); + assert (acc_is_present (block2, SIZE)); + assert (!acc_is_present (block3, SIZE)); + +#ifdef OPENACC_API + acc_copyout (block2, SIZE); +#else +#pragma acc exit data copyout(block2[0:SIZE]) +#endif + assert (!acc_is_present (block2, SIZE)); + + free (block1); + free (block2); + free (block3); +} + +int +main (int argc, char *argv[]) +{ + f1 (); + f2 (); + f3 (); + f4 (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c new file mode 100644 index 00000000000..8507a0586a5 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c @@ -0,0 +1,5 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "static-dynamic-lifetimes-6.c" +/* { dg-output "libgomp: cannot handle delete/copyout within data region" } */ +/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c new file mode 100644 index 00000000000..ca3b385fbcc --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c @@ -0,0 +1,46 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#define SIZE 1024 + +int +main (int argc, char *argv[]) +{ + char *block1 = (char *) malloc (SIZE); + char *block2 = (char *) malloc (SIZE); + +#ifdef OPENACC_API + acc_copyin (block1, SIZE); + acc_copyin (block2, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE], block2[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyout (block1, SIZE); + acc_copyout (block2, SIZE); + /* Error output checked in static-dynamic-lifetimes-6-lib.c. */ +#else +#pragma acc exit data copyout(block1[0:SIZE], block2[0:SIZE]) +/* We can only do this for a single dynamic data mapping at present. */ +/* { dg-output "libgomp: cannot handle .exit data. within data region" } */ +/* { dg-shouldfail "" } */ +#endif + /* These should stay present until the end of the static data lifetime. */ + assert (acc_is_present (block1, SIZE)); + assert (acc_is_present (block2, SIZE)); + } + + assert (!acc_is_present (block1, SIZE)); + assert (!acc_is_present (block2, SIZE)); + + free (block1); + free (block2); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c new file mode 100644 index 00000000000..962b5926f79 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c @@ -0,0 +1,5 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "static-dynamic-lifetimes-7.c" +/* { dg-output "libgomp: cannot handle delete/copyout within data region" } */ +/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c new file mode 100644 index 00000000000..dfcc7cae961 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c @@ -0,0 +1,45 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#define SIZE 1024 + +int +main (int argc, char *argv[]) +{ + char *block1 = (char *) malloc (SIZE); + char *block2 = (char *) malloc (SIZE); + +#ifdef OPENACC_API + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE]) + { +/* We can't attach the dynamic data mapping's (block1) target_mem_desc to the + enclosing static data region here, because that region maps block2 also. */ +#ifdef OPENACC_API + acc_copyout (block1, SIZE); + /* Error output checked in static-dynamic-lifetimes-7-lib.c. */ +#else +#pragma acc exit data copyout(block1[0:SIZE]) +/* { dg-output "libgomp: cannot handle .exit data. within data region" } */ +/* { dg-shouldfail "" } */ +#endif + /* These should stay present until the end of the static data lifetime. */ + assert (acc_is_present (block1, SIZE)); + assert (acc_is_present (block2, SIZE)); + } + + assert (!acc_is_present (block1, SIZE)); + assert (!acc_is_present (block2, SIZE)); + + free (block1); + free (block2); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c new file mode 100644 index 00000000000..2581d7e2559 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c @@ -0,0 +1,5 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ +/* { dg-additional-options "-DOPENACC_API" } */ +#include "static-dynamic-lifetimes-8.c" +/* { dg-output "libgomp: cannot handle delete/copyout within data region" } */ +/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c new file mode 100644 index 00000000000..e3a64399fe9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c @@ -0,0 +1,50 @@ +/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */ + +#include +#include +#include + +#define SIZE 1024 + +int +main (int argc, char *argv[]) +{ + char *block1 = (char *) malloc (SIZE); + char *block2 = (char *) malloc (SIZE); + +#ifdef OPENACC_API + acc_copyin (block1, SIZE); +#else +#pragma acc enter data copyin(block1[0:SIZE]) +#endif + +#pragma acc data copy(block1[0:SIZE], block2[0:SIZE]) + { +#ifdef OPENACC_API + acc_copyout (block1, SIZE); + acc_copyin (block2, SIZE); + /* Error output checked in static-dynamic-lifetimes-8-lib.c. */ +#else +#pragma acc exit data copyout(block1[0:SIZE]) +/* { dg-output "libgomp: cannot handle .exit data. within data region" } */ +/* { dg-shouldfail "" } */ +#pragma acc enter data copyin(block2[0:SIZE]) +#endif + assert (acc_is_present (block1, SIZE)); + assert (acc_is_present (block2, SIZE)); + } + + assert (!acc_is_present (block1, SIZE)); + assert (acc_is_present (block2, SIZE)); +#ifdef OPENACC_API + acc_copyout (block2, SIZE); +#else +#pragma acc exit data copyout(block2[0:SIZE]) +#endif + assert (!acc_is_present (block2, SIZE)); + + free (block1); + free (block2); + + return 0; +}