From patchwork Thu Jun 22 07:18:36 2017 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 779363 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org 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 3wtXyZ3gPFz9s82 for ; Thu, 22 Jun 2017 17:19:33 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="hO5PqfCQ"; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:to:cc:references:from:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=T5qd8MShyErrXEfIM kCmvdfkzLTeU8o7BOkZqq9/3xqiHpTlNVxdKRzZhBS+nsuV+mqs1oz681vzHpeO1 EuOiXac1jqxqsRrbaI8DjNvXSaTFXyJBBek6h69hg+16mFQt0qyrjXers6xQ8ZLN SUgZ7U852/6Q535glJ6Azr1sd0= 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 :subject:to:cc:references:from:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=ZzIETGioDUnnc6X9vaEYLsE /vkw=; b=hO5PqfCQ/23WAkh24ndx465V98FBcS+d4Ko/U6C9yV/lrkO0ElIhL3U p57+6v/ICZ0XlpeRrdM792kj282QvS5QsMPpay3K7WopPtlSeFOlKgW4aEURNPBq GomedZG6l4pHkK/GPPPSCID1a1Fxq5WgV2UYVa9zDN7UrXJBd2ak= Received: (qmail 122872 invoked by alias); 22 Jun 2017 07:19:22 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 122857 invoked by uid 89); 22 Jun 2017 07:19:19 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-10.2 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_2, GIT_PATCH_3, KAM_ASCII_DIVIDERS, RCVD_IN_DNSWL_NONE, SPF_PASS, URIBL_RED autolearn=ham version=3.3.2 spammy=ji, forgotten, Cesar, cesar X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Thu, 22 Jun 2017 07:19:16 +0000 Received: from svr-orw-mbx-06.mgc.mentorg.com ([147.34.90.206]) by relay1.mentorg.com with esmtp id 1dNwOW-0000WD-VH from ChungLin_Tang@mentor.com ; Thu, 22 Jun 2017 00:19:12 -0700 Received: from svr-orw-mbx-08.mgc.mentorg.com (147.34.90.208) by SVR-ORW-MBX-06.mgc.mentorg.com (147.34.90.206) with Microsoft SMTP Server (TLS) id 15.0.1263.5; Thu, 22 Jun 2017 00:19:10 -0700 Received: from [0.0.0.0] (147.34.91.1) by svr-orw-mbx-08.mgc.mentorg.com (147.34.90.208) with Microsoft SMTP Server (TLS) id 15.0.1263.5 via Frontend Transport; Thu, 22 Jun 2017 00:19:08 -0700 Subject: Re: [libgomp, OpenACC] Add more map handling for enter/exit data directives To: Jakub Jelinek CC: gcc-patches , Thomas Schwinge , Cesar Philippidis References: <8f8391b0-5d9a-19db-671f-81f16063da59@mentor.com> <20170613160025.GK2099@tucnak> From: Chung-Lin Tang Message-ID: <02d54864-7326-d44e-088b-bda2f0b3a136@mentor.com> Date: Thu, 22 Jun 2017 15:18:36 +0800 User-Agent: Mozilla/5.0 (Windows NT 6.1; WOW64; rv:52.0) Gecko/20100101 Thunderbird/52.2.0 MIME-Version: 1.0 In-Reply-To: <20170613160025.GK2099@tucnak> On 2017/6/14 12:00 AM, Jakub Jelinek wrote: > I don't see sufficient information on what you want to change and why > and whether the changes are backwards compatible (say will a valid > OpenACC 2.0 program compiled by GCC 7 work against both libgomp from GCC 7 > as well as one with this patch)? > Can you write a few paragraphs on it (doesn't have to be comments in the > source, mailing list is fine)? The current code doesn't handle GOMP_MAP_TO (present_or_copyin) and also the GOMP_MAP_PSET/MAP_POINTER handling wasn't entirely correct. This patch fixes them. In the new attached patch, I added a fix a memory management fix that was forgotten earlier. The collective patch was originally by Cesar, from here: https://gcc.gnu.org/ml/gcc-patches/2015-05/msg01367.html (in that post, his main emphasis was the memory management fix) >> + case GOMP_MAP_ALLOC: >> + acc_present_or_create (hostaddrs[i], sizes[i]); >> break; >> case GOMP_MAP_FORCE_ALLOC: >> acc_create (hostaddrs[i], sizes[i]); >> break; >> - case GOMP_MAP_FORCE_PRESENT: >> + case GOMP_MAP_TO: >> acc_present_or_copyin (hostaddrs[i], sizes[i]); >> break; >> case GOMP_MAP_FORCE_TO: >> - acc_present_or_copyin (hostaddrs[i], sizes[i]); >> + acc_copyin (hostaddrs[i], sizes[i]); >> break; > > E.g. in this hunk you remove GOMP_MAP_POINTER and GOMP_MAP_FORCE_PRESENT > handling and significantly change GOMP_MAP_FORCE_TO. The first two will > now gomp_fatal, right? Can it ever appear in GOACC_enter_exit_data > calls? GOMP_MAP_FORCE_PRESENT does not appear in enter/exit data directives, while GOMP_MAP_POINTER is handled in find_pointer(). Thanks, Chung-Lin 2017-06-22 Cesar Philippidis Thomas Schwinge Chung-Lin Tang libgomp/ * oacc-mem.c (gomp_acc_remove_pointer): Fix a memory leak preventing target_mem_desc.to_free from being deallocated with acc exit data. * oacc-parallel.c (find_pset): Adjust and rename from... (find_pointer): ...this function. (GOACC_enter_exit_data): Handle GOMP_MAP_TO and GOMP_MAP_ALLOC, adjust find_pointer calls into find_pset, adjust pointer map handling, add acc_is_present guards to calls to gomp_acc_insert_pointer and gomp_acc_remove_pointer. * testsuite/libgomp.oacc-c-c++-common/data-2.c: Update test. * testsuite/libgomp.oacc-c-c++-common/enter-data.c: New test. * testsuite/libgomp.oacc-fortran/data-2.f90: Update test. Index: oacc-mem.c =================================================================== --- oacc-mem.c (revision 249091) +++ oacc-mem.c (working copy) @@ -698,10 +698,8 @@ gomp_acc_remove_pointer (void *h, bool force_copyf if (t->refcount == minrefs) { /* This is the last reference, so pull the descriptor off the - chain. This avoids gomp_unmap_vars via gomp_unmap_tgt from + chain. This prevents gomp_unmap_vars via gomp_unmap_tgt from freeing the device memory. */ - t->tgt_end = 0; - t->to_free = 0; for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL; tp = t, t = t->prev) @@ -717,8 +715,7 @@ gomp_acc_remove_pointer (void *h, bool force_copyf } } - if (force_copyfrom) - t->list[0].copy_from = 1; + t->list[0]->copy_from = force_copyfrom ? 1 : 0; gomp_mutex_unlock (&acc_dev->lock); Index: oacc-parallel.c =================================================================== --- oacc-parallel.c (revision 249091) +++ oacc-parallel.c (working copy) @@ -38,15 +38,23 @@ #include #include +/* Returns the number of mappings associated with the pointer or pset. PSET + have three mappings, whereas pointer have two. */ + static int -find_pset (int pos, size_t mapnum, unsigned short *kinds) +find_pointer (int pos, size_t mapnum, unsigned short *kinds) { if (pos + 1 >= mapnum) return 0; unsigned char kind = kinds[pos+1] & 0xff; - return kind == GOMP_MAP_TO_PSET; + if (kind == GOMP_MAP_TO_PSET) + return 3; + else if (kind == GOMP_MAP_POINTER) + return 2; + + return 0; } static void goacc_wait (int async, int num_waits, va_list *ap); @@ -298,7 +306,9 @@ GOACC_enter_exit_data (int device, size_t mapnum, if (kind == GOMP_MAP_FORCE_ALLOC || kind == GOMP_MAP_FORCE_PRESENT - || kind == GOMP_MAP_FORCE_TO) + || kind == GOMP_MAP_FORCE_TO + || kind == GOMP_MAP_TO + || kind == GOMP_MAP_ALLOC) { data_enter = true; break; @@ -312,31 +322,39 @@ GOACC_enter_exit_data (int device, size_t mapnum, kind); } + /* In c, non-pointers and arrays are represented by a single data clause. + Dynamically allocated arrays and subarrays are represented by a data + clause followed by an internal GOMP_MAP_POINTER. + + In fortran, scalars and not allocated arrays are represented by a + single data clause. Allocated arrays and subarrays have three mappings: + 1) the original data clause, 2) a PSET 3) a pointer to the array data. + */ + if (data_enter) { for (i = 0; i < mapnum; i++) { unsigned char kind = kinds[i] & 0xff; - /* Scan for PSETs. */ - int psets = find_pset (i, mapnum, kinds); + /* Scan for pointers and PSETs. */ + int pointer = find_pointer (i, mapnum, kinds); - if (!psets) + if (!pointer) { switch (kind) { - case GOMP_MAP_POINTER: - gomp_acc_insert_pointer (1, &hostaddrs[i], &sizes[i], - &kinds[i]); + case GOMP_MAP_ALLOC: + acc_present_or_create (hostaddrs[i], sizes[i]); break; case GOMP_MAP_FORCE_ALLOC: acc_create (hostaddrs[i], sizes[i]); break; - case GOMP_MAP_FORCE_PRESENT: + case GOMP_MAP_TO: acc_present_or_copyin (hostaddrs[i], sizes[i]); break; case GOMP_MAP_FORCE_TO: - acc_present_or_copyin (hostaddrs[i], sizes[i]); + acc_copyin (hostaddrs[i], sizes[i]); break; default: gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", @@ -346,12 +364,16 @@ GOACC_enter_exit_data (int device, size_t mapnum, } else { - gomp_acc_insert_pointer (3, &hostaddrs[i], &sizes[i], &kinds[i]); + if (!acc_is_present (hostaddrs[i], sizes[i])) + { + gomp_acc_insert_pointer (pointer, &hostaddrs[i], + &sizes[i], &kinds[i]); + } /* Increment 'i' by two because OpenACC requires fortran arrays to be contiguous, so each PSET is associated with one of MAP_FORCE_ALLOC/MAP_FORCE_PRESET/MAP_FORCE_TO, and one MAP_POINTER. */ - i += 2; + i += pointer - 1; } } } @@ -360,19 +382,15 @@ GOACC_enter_exit_data (int device, size_t mapnum, { unsigned char kind = kinds[i] & 0xff; - int psets = find_pset (i, mapnum, kinds); + int pointer = find_pointer (i, mapnum, kinds); - if (!psets) + if (!pointer) { switch (kind) { - case GOMP_MAP_POINTER: - gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff) - == GOMP_MAP_FORCE_FROM, - async, 1); - break; case GOMP_MAP_DELETE: - acc_delete (hostaddrs[i], sizes[i]); + if (acc_is_present (hostaddrs[i], sizes[i])) + acc_delete (hostaddrs[i], sizes[i]); break; case GOMP_MAP_FORCE_FROM: acc_copyout (hostaddrs[i], sizes[i]); @@ -385,10 +403,14 @@ GOACC_enter_exit_data (int device, size_t mapnum, } else { - gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff) - == GOMP_MAP_FORCE_FROM, async, 3); - /* See the above comment. */ - i += 2; + if (acc_is_present (hostaddrs[i], sizes[i])) + { + gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff) + == GOMP_MAP_FORCE_FROM, async, + pointer); + /* See the above comment. */ + } + i += pointer - 1; } } Index: testsuite/libgomp.oacc-c-c++-common/data-2.c =================================================================== --- testsuite/libgomp.oacc-c-c++-common/data-2.c (revision 249091) +++ testsuite/libgomp.oacc-c-c++-common/data-2.c (working copy) @@ -3,6 +3,7 @@ /* { dg-do run } */ #include +#include int main (int argc, char **argv) @@ -32,7 +33,7 @@ main (int argc, char **argv) for (i = 0; i < N; i++) b[i] = a[i]; -#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait async +#pragma acc exit data copyout (a[0:N], b[0:N]) delete (N) wait async #pragma acc wait for (i = 0; i < N; i++) @@ -46,6 +47,32 @@ main (int argc, char **argv) for (i = 0; i < N; i++) { + a[i] = 3.0; + b[i] = 0.0; + } + +#pragma acc enter data copyin (a[0:N]) async +#pragma acc enter data copyin (b[0:N]) async wait +#pragma acc enter data copyin (N) async wait +#pragma acc parallel async wait present (a[0:N]) present (b[0:N]) present (N) +#pragma acc loop + for (i = 0; i < N; i++) + b[i] = a[i]; + +#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) delete (N) wait async +#pragma acc wait + + for (i = 0; i < N; i++) + { + if (a[i] != 3.0) + abort (); + + if (b[i] != 3.0) + abort (); + } + + for (i = 0; i < N; i++) + { a[i] = 2.0; b[i] = 0.0; } @@ -56,7 +83,7 @@ main (int argc, char **argv) for (i = 0; i < N; i++) b[i] = a[i]; -#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait (1) async (1) +#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) delete (N) wait (1) async (1) #pragma acc wait (1) for (i = 0; i < N; i++) @@ -93,7 +120,7 @@ main (int argc, char **argv) for (i = 0; i < N; i++) d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; -#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) wait (1, 2, 3) async (1) +#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) delete (N) wait (1, 2, 3) async (1) #pragma acc wait (1) for (i = 0; i < N; i++) @@ -161,5 +188,156 @@ main (int argc, char **argv) abort (); } +#if !ACC_MEM_SHARED + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + } + +#pragma acc enter data present_or_copyin (a[0:N]) + + if (!acc_is_present (a, nbytes)) + abort (); + +#pragma acc exit data copyout (a[0:N]) + + if (acc_is_present (a, nbytes)) + abort (); + +#pragma acc exit data delete (a[0:N]) + + if (acc_is_present (a, nbytes)) + abort (); + +#pragma acc enter data present_or_copyin (a[0:N], b[0:N]) + + if (!acc_is_present (a, nbytes)) + abort (); + + if (!acc_is_present (b, nbytes)) + abort (); + +#pragma acc exit data copyout (a[0:N], b[0:N]) + + if (acc_is_present (a, nbytes)) + abort (); + + if (acc_is_present (b, nbytes)) + abort (); + +#pragma acc exit data delete (a[0:N], b[0:N]) + + if (acc_is_present (a, nbytes)) + abort (); + + if (acc_is_present (b, nbytes)) + abort (); + +#pragma acc enter data present_or_create (a[0:N], b[0:N]) + + if (!acc_is_present (a, nbytes)) + abort (); + + if (!acc_is_present (b, nbytes)) + abort (); + +#pragma acc exit data copyout (a[0:N], b[0:N]) + + if (acc_is_present (a, nbytes)) + abort (); + + if (acc_is_present (b, nbytes)) + abort (); + +#pragma acc enter data present_or_create (a[0:N], b[0:N]) + + if (!acc_is_present (a, nbytes)) + abort (); + + if (!acc_is_present (b, nbytes)) + abort (); + +#pragma acc enter data present_or_create (a[0:N], b[0:N]) + + if (!acc_is_present (a, nbytes)) + abort (); + + if (!acc_is_present (b, nbytes)) + abort (); + +#pragma acc exit data delete (a[0:N], b[0:N]) + + if (acc_is_present (a, nbytes)) + abort (); + + if (acc_is_present (b, nbytes)) + abort (); + +#pragma acc exit data delete (a[0:N], b[0:N]) + + if (acc_is_present (a, nbytes)) + abort (); + + if (acc_is_present (b, nbytes)) + abort (); + +#pragma acc enter data create (a[0:N], b[0:N]) + + if (!acc_is_present (a, nbytes)) + abort (); + + if (!acc_is_present (b, nbytes)) + abort (); + +#pragma acc enter data present_or_copyin (a[0:N], b[0:N]) + + if (!acc_is_present (a, nbytes)) + abort (); + + if (!acc_is_present (b, nbytes)) + abort (); + +#pragma acc exit data delete (a[0:N], b[0:N]) + + if (acc_is_present (a, nbytes)) + abort (); + + if (acc_is_present (b, nbytes)) + abort (); + +#pragma acc exit data delete (a[0:N], b[0:N]) + + if (acc_is_present (a, nbytes)) + abort (); + + if (acc_is_present (b, nbytes)) + abort (); + +#pragma acc enter data present_or_copyin (a[0:N]) + + if (!acc_is_present (a, nbytes)) + abort (); + + if (acc_is_present (b, nbytes)) + abort (); + +#pragma acc exit data delete (a[0:N], b[0:N]) + + if (acc_is_present (a, nbytes)) + abort (); + + if (acc_is_present (b, nbytes)) + abort (); + +#pragma acc exit data delete (a[0:N], b[0:N]) + + if (acc_is_present (a, nbytes)) + abort (); + + if (acc_is_present (b, nbytes)) + abort (); +#endif + return 0; } Index: testsuite/libgomp.oacc-c-c++-common/enter-data.c =================================================================== --- testsuite/libgomp.oacc-c-c++-common/enter-data.c (revision 0) +++ testsuite/libgomp.oacc-c-c++-common/enter-data.c (revision 0) @@ -0,0 +1,23 @@ +/* This test verifies that the present data clauses to acc enter data + don't cause duplicate mapping failures at runtime. */ + +/* { dg-do run } */ + +#include + +int +main (void) +{ + int a; + +#pragma acc enter data copyin (a) +#pragma acc enter data pcopyin (a) +#pragma acc enter data pcreate (a) +#pragma acc exit data delete (a) + +#pragma acc enter data create (a) +#pragma acc enter data pcreate (a) +#pragma acc exit data delete (a) + + return 0; +} Index: testsuite/libgomp.oacc-fortran/data-2.f90 =================================================================== --- testsuite/libgomp.oacc-fortran/data-2.f90 (revision 249091) +++ testsuite/libgomp.oacc-fortran/data-2.f90 (working copy) @@ -1,9 +1,16 @@ ! { dg-do run } +! { dg-additional-options "-cpp" } program test + use openacc integer, parameter :: N = 8 real, allocatable :: a(:,:), b(:,:) + real, allocatable :: c(:), d(:) + integer i, j + i = 0 + j = 0 + allocate (a(N,N)) allocate (b(N,N)) @@ -12,7 +19,7 @@ program test !$acc enter data copyin (a(1:N,1:N), b(1:N,1:N)) - !$acc parallel + !$acc parallel present (a(1:N,1:N), b(1:N,1:N)) do i = 1, n do j = 1, n b(j,i) = a (j,i) @@ -28,4 +35,171 @@ program test if (b(j,i) .ne. 3.0) call abort end do end do + + allocate (c(N)) + allocate (d(N)) + + c(:) = 3.0 + d(:) = 0.0 + + !$acc enter data copyin (c(1:N)) create (d(1:N)) async + !$acc wait + + !$acc parallel present (c(1:N), d(1:N)) + do i = 1, N + d(i) = c(i) + 1 + end do + !$acc end parallel + + !$acc exit data copyout (c(1:N), d(1:N)) async + !$acc wait + + do i = 1, N + if (d(i) .ne. 4.0) call abort + end do + + c(:) = 3.0 + d(:) = 0.0 + + !$acc enter data copyin (c(1:N)) async + !$acc enter data create (d(1:N)) wait + !$acc wait + + !$acc parallel present (c(1:N), d(1:N)) + do i = 1, N + d(i) = c(i) + 1 + end do + !$acc end parallel + + !$acc exit data delete (c(1:N)) copyout (d(1:N)) async + !$acc exit data async + !$acc wait + + do i = 1, N + if (d(i) .ne. 4.0) call abort + end do + +#if !ACC_MEM_SHARED + + c(:) = 3.0 + d(:) = 0.0 + + !$acc enter data present_or_copyin (c(0:N)) + + if (acc_is_present (c) .eqv. .FALSE.) call abort + + !$acc exit data copyout (c(0:N)) + + if (acc_is_present (c) .eqv. .TRUE.) call abort + + !$acc exit data delete (c(0:N)) + + if (acc_is_present (c) .eqv. .TRUE.) call abort + + do i = 1, N + if (c(i) .ne. 3.0) call abort + end do + + c(:) = 5.0 + d(:) = 9.0 + + !$acc enter data present_or_copyin (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .FALSE.) call abort + if (acc_is_present (d) .eqv. .FALSE.) call abort + + !$acc exit data copyout (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .TRUE.) call abort + if (acc_is_present (d) .eqv. .TRUE.) call abort + + !$acc exit data delete (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .TRUE.) call abort + if (acc_is_present (d) .eqv. .TRUE.) call abort + + do i = 1, N + if (c(i) .ne. 5.0) call abort + if (d(i) .ne. 9.0) call abort + end do + + !$acc enter data present_or_create (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .FALSE.) call abort + if (acc_is_present (d) .eqv. .FALSE.) call abort + + !$acc parallel present (c(0:N), d(0:N)) + do i = 1, N + c(i) = 1.0; + d(i) = 2.0; + end do + !$acc end parallel + + !$acc exit data copyout (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .TRUE.) call abort + if (acc_is_present (d) .eqv. .TRUE.) call abort + + do i = 1, N + if (c(i) .ne. 1.0) call abort + if (d(i) .ne. 2.0) call abort + end do + + !$acc enter data present_or_create (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .FALSE.) call abort + if (acc_is_present (d) .eqv. .FALSE.) call abort + + !$acc enter data present_or_create (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .FALSE.) call abort + if (acc_is_present (d) .eqv. .FALSE.) call abort + + !$acc exit data delete (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .TRUE.) call abort + if (acc_is_present (d) .eqv. .TRUE.) call abort + + !$acc exit data delete (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .TRUE.) call abort + if (acc_is_present (d) .eqv. .TRUE.) call abort + + !$acc enter data create (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .FALSE.) call abort + if (acc_is_present (d) .eqv. .FALSE.) call abort + + !$acc enter data present_or_copyin (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .FALSE.) call abort + if (acc_is_present (d) .eqv. .FALSE.) call abort + + !$acc exit data delete (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .TRUE.) call abort + if (acc_is_present (d) .eqv. .TRUE.) call abort + + !$acc exit data delete (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .TRUE.) call abort + if (acc_is_present (d) .eqv. .TRUE.) call abort + + !$acc enter data present_or_copyin (c(0:N)) + + if (acc_is_present (c) .eqv. .FALSE.) call abort + if (acc_is_present (d) .eqv. .TRUE.) call abort + + !$acc exit data delete (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .TRUE.) call abort + if (acc_is_present (d) .eqv. .TRUE.) call abort + + !$acc exit data delete (c(0:N), d(0:N)) + + if (acc_is_present (c) .eqv. .TRUE.) call abort + if (acc_is_present (d) .eqv. .TRUE.) call abort + +#endif + end program test