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; +}