From patchwork Wed Oct 25 09:29:52 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1854936 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4SFkDd6HkMz23jV for ; Wed, 25 Oct 2023 20:30:25 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id E41203857345 for ; Wed, 25 Oct 2023 09:30:23 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 6C0103858281; Wed, 25 Oct 2023 09:30:04 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 6C0103858281 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 6C0103858281 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=68.232.141.98 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1698226210; cv=none; b=suG7gzboO1EYzoNwALwe8jwfkBm2M+Er75H1rsfeuMblakyTBRx5tmN85JlIzswbRfV3h/RazI19Ay0gtLHPFllImY21dlhRitnvspSZJ/cN2wPUZ8y4M3s5/5ghVNW0EPeFYNNTMI9YtlfAY9EwBn2XvVxuQ0Nl8eEua89Xrvk= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1698226210; c=relaxed/simple; bh=jXThVVH3UEVIt44B73pzODJRCSoGXxDJ5uKKmJ/jlBg=; h=From:To:Subject:Date:Message-ID:MIME-Version; b=irOfvcoHkmfmj5kN0wXIEjMVUNQhpSM9zpQBJE20KcWXEpgGG2BgM9TUeVwOrw5VY/qDli/d+rc/qye1Es2i6TGjDtOiU1LzJDnw1v6nutsMvP7fAcsnEUHxvVmmOSCOVj0802pPmVd2m1lD4YtJnxi42jJ3kYaLGqTo2FSBHkg= ARC-Authentication-Results: i=1; server2.sourceware.org X-CSE-ConnectionGUID: Glf/cX+FSjqBUx1wZ6CfkQ== X-CSE-MsgGUID: BgwEOiVjTlGlI3vvICKC5A== X-IronPort-AV: E=Sophos;i="6.03,250,1694764800"; d="scan'208,223";a="23200093" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 25 Oct 2023 01:30:03 -0800 IronPort-SDR: q8GaMWBk8gpKXCteScL+8FTV9RFRoC8l4nFXHro7ATknhq0wCW15VsYFW56/V1NvmFPjT1WDXz oR0j9nXGnX9UMshttjopL7lDpPb0aNz2gcbgYViphGDGB5y009Vt7CVncjJdfqO7ttuyZNQAZi x7fBuzUMi0GElf9Ge0C6r3lNICkDY7bRYPm372HBY3/GNAxPxtkAOmA7M6b5mrdi8DwizkX8UF 3dQ2cBynFxvpOjviSSyT4o5/Q64xzFvbVirgzFOWBjLgHcdsHvdvhoRfbQLslIbW3farQ3N0wu M+w= From: Thomas Schwinge To: Chung-Lin Tang , , CC: Catherine Moore , Tobias Burnus Subject: Extend test suite coverage for OpenACC 'self' clause for compute constructs (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs) In-Reply-To: <87pm13w04d.fsf@euler.schwinge.homeip.net> References: <87pm13w04d.fsf@euler.schwinge.homeip.net> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/28.2 (x86_64-pc-linux-gnu) Date: Wed, 25 Oct 2023 11:29:52 +0200 Message-ID: <87edhjvylr.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) To svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) X-Spam-Status: No, score=-11.7 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_LOTSOFHASH, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Hi! On 2023-10-25T10:57:06+0200, I wrote: > With minor textual conflicts resolved, I've pushed this to master branch > in commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a > "OpenACC 2.7: Implement self clause for compute constructs", see > attached. > > > I'll then apply/submit a number of follow-on commits. > From 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a Mon Sep 17 00:00:00 2001 > From: Chung-Lin Tang > Date: Tue, 13 Jun 2023 08:44:31 -0700 > Subject: [PATCH] OpenACC 2.7: Implement self clause for compute constructs > .../c-c++-common/goacc/self-clause-1.c | 22 + > .../c-c++-common/goacc/self-clause-2.c | 17 + > gcc/testsuite/gfortran.dg/goacc/self.f95 | 53 + > .../libgomp.oacc-c-c++-common/self-1.c | 962 ++++++++++++++++++ I found that insufficient, and added some more. Pushed to master branch commit 047841a68ebf5f991e842961f9e54f3c10b94f2c "Extend test suite coverage for OpenACC 'self' clause for compute constructs", see attached. This is mostly just adapting and cross-linking some existing 'if' clause test cases. (..., which turned up a problem when the 'self' clause is used with OpenACC 'kernels'.) Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 From 047841a68ebf5f991e842961f9e54f3c10b94f2c Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Mon, 23 Oct 2023 14:53:29 +0200 Subject: [PATCH] Extend test suite coverage for OpenACC 'self' clause for compute constructs ... on top of what was provided in recent commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a "OpenACC 2.7: Implement self clause for compute constructs". gcc/testsuite/ * c-c++-common/goacc/if-clause-2.c: Enhance. * c-c++-common/goacc/self-clause-1.c: Likewise. * c-c++-common/goacc/self-clause-2.c: Likewise. * gfortran.dg/goacc/if.f95: Likewise. * gfortran.dg/goacc/kernels-tree.f95: Likewise. * gfortran.dg/goacc/parallel-tree.f95: Likewise. * gfortran.dg/goacc/self.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/if-1.c: Enhance. * testsuite/libgomp.oacc-c-c++-common/self-1.c: Likewise. * testsuite/libgomp.oacc-fortran/if-1.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/if-self-1.c: New. * testsuite/libgomp.oacc-fortran/self-1.f90: Likewise. --- .../c-c++-common/goacc/if-clause-2.c | 2 + .../c-c++-common/goacc/self-clause-1.c | 6 + .../c-c++-common/goacc/self-clause-2.c | 20 + gcc/testsuite/gfortran.dg/goacc/if.f95 | 10 +- .../gfortran.dg/goacc/kernels-tree.f95 | 5 +- .../gfortran.dg/goacc/parallel-tree.f95 | 3 +- gcc/testsuite/gfortran.dg/goacc/self.f95 | 8 + .../libgomp.oacc-c-c++-common/if-1.c | 4 + .../libgomp.oacc-c-c++-common/if-self-1.c | 36 + .../libgomp.oacc-c-c++-common/self-1.c | 5 + .../testsuite/libgomp.oacc-fortran/if-1.f90 | 4 + .../testsuite/libgomp.oacc-fortran/self-1.f90 | 996 ++++++++++++++++++ 12 files changed, 1094 insertions(+), 5 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 diff --git a/gcc/testsuite/c-c++-common/goacc/if-clause-2.c b/gcc/testsuite/c-c++-common/goacc/if-clause-2.c index a48072509e1..71475521758 100644 --- a/gcc/testsuite/c-c++-common/goacc/if-clause-2.c +++ b/gcc/testsuite/c-c++-common/goacc/if-clause-2.c @@ -1,3 +1,5 @@ +/* See also 'self-clause-2.c'. */ + /* { dg-additional-options "-fdump-tree-gimple" } */ /* { dg-additional-options "--param=openacc-kernels=decompose" } { dg-additional-options "-fdump-tree-omp_oacc_kernels_decompose" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-1.c b/gcc/testsuite/c-c++-common/goacc/self-clause-1.c index fe892bea210..28de3dc0584 100644 --- a/gcc/testsuite/c-c++-common/goacc/self-clause-1.c +++ b/gcc/testsuite/c-c++-common/goacc/self-clause-1.c @@ -5,6 +5,8 @@ f (int b) { struct { int i; } *p; +#pragma acc parallel self(0) self(b) /* { dg-error "too many 'self' clauses" } */ + ; #pragma acc parallel self self(b) /* { dg-error "too many 'self' clauses" } */ ; #pragma acc parallel self(*p) @@ -12,6 +14,8 @@ f (int b) { dg-error {could not convert '\* p' from 'f\(int\)::' to 'bool'} {} { target c++ } .-2 } */ ; +#pragma acc kernels self(0) self(b) /* { dg-error "too many 'self' clauses" } */ + ; #pragma acc kernels self self(b) /* { dg-error "too many 'self' clauses" } */ ; #pragma acc kernels self(*p) @@ -19,6 +23,8 @@ f (int b) { dg-error {could not convert '\* p' from 'f\(int\)::' to 'bool'} {} { target c++ } .-2 } */ ; +#pragma acc serial self(0) self(b) /* { dg-error "too many 'self' clauses" } */ + ; #pragma acc serial self self(b) /* { dg-error "too many 'self' clauses" } */ ; #pragma acc serial self(*p) diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c index d932ac9a4a6..769694baec9 100644 --- a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c +++ b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c @@ -1,3 +1,5 @@ +/* See also 'if-clause-2.c'. */ + /* { dg-additional-options "-fdump-tree-gimple" } */ void @@ -15,3 +17,21 @@ f (short c) /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_serial map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */ ++c; } + +/* The same, but with implicit 'true' condition-argument. */ + +void +g (short d) +{ +#pragma acc parallel self copy(d) + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */ + ++d; + +#pragma acc kernels self copy(d) + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_kernels map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */ + ++d; + +#pragma acc serial self copy(d) + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_serial map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */ + ++d; +} diff --git a/gcc/testsuite/gfortran.dg/goacc/if.f95 b/gcc/testsuite/gfortran.dg/goacc/if.f95 index 56f3711f320..753ef8251c2 100644 --- a/gcc/testsuite/gfortran.dg/goacc/if.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/if.f95 @@ -1,3 +1,5 @@ +! See also 'self.f95'. + ! { dg-do compile } program test @@ -12,12 +14,14 @@ program test !$acc end parallel !$acc parallel if (1) ! { dg-error "scalar LOGICAL expression" } !$acc end parallel - !$acc kernels if (i) ! { dg-error "scalar LOGICAL expression" } - !$acc end kernels + !$acc kernels if ! { dg-error "Expected '\\(' after 'if'" } !$acc kernels if () ! { dg-error "Invalid character" } + !$acc kernels if (i) ! { dg-error "scalar LOGICAL expression" } + !$acc end kernels !$acc kernels if (1) ! { dg-error "scalar LOGICAL expression" } !$acc end kernels + !$acc data if ! { dg-error "Expected '\\(' after 'if'" } !$acc data if () ! { dg-error "Invalid character" } !$acc data if (i) ! { dg-error "scalar LOGICAL expression" } @@ -36,12 +40,14 @@ program test !$acc end parallel !$acc parallel if (i.gt.1) !$acc end parallel + !$acc kernels if (x) !$acc end kernels !$acc kernels if (.true.) !$acc end kernels !$acc kernels if (i.gt.1) !$acc end kernels + !$acc data if (x) !$acc end data !$acc data if (.true.) diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 index ceb07fbb9e9..1ba04a84e12 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 @@ -12,6 +12,7 @@ program test logical :: l = .true. !$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) & + !$acc self & !$acc copy(i), copyin(j), copyout(k), create(m) & !$acc no_create(n) & !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) & @@ -27,7 +28,7 @@ end program test ! { dg-final { scan-tree-dump-times "num_gangs" 1 "original" } } ! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } } ! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } - +! { dg-final { scan-tree-dump-times "self\\(1\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } } @@ -42,4 +43,4 @@ end program test ! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels if\((?:D\.|_)[0-9]+\)$} 1 "omp_oacc_kernels_decompose" } } -! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single num_gangs\(1\) if\((?:D\.|_)[0-9]+\) async\(-1\)$} 1 "omp_oacc_kernels_decompose" } } +! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single num_gangs\(1\) if\((?:D\.|_)[0-9]+\) self\(1\) async\(-1\)$} 1 "omp_oacc_kernels_decompose" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 index 6110d93b91e..0d4ec1133af 100644 --- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 @@ -14,6 +14,7 @@ program test logical :: l = .true. !$acc parallel if(l) async num_gangs(i) num_workers(i) vector_length(i) & + !$acc self & !$acc reduction(max:q), copy(i), copyin(j), copyout(k), create(m) & !$acc no_create(n) & !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) & @@ -33,7 +34,7 @@ end program test ! { dg-final { scan-tree-dump-times "num_gangs" 1 "original" } } ! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } } ! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } - +! { dg-final { scan-tree-dump-times "self\\(1\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "reduction\\(max:q\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/self.f95 b/gcc/testsuite/gfortran.dg/goacc/self.f95 index 4817f16be56..aa0f6fe88c5 100644 --- a/gcc/testsuite/gfortran.dg/goacc/self.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/self.f95 @@ -1,3 +1,5 @@ +! See also 'if.f95'. + ! { dg-do compile } program test @@ -29,6 +31,8 @@ program test !$acc kernels self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" } !$acc serial self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" } + !$acc parallel self + !$acc end parallel !$acc parallel self (x) !$acc end parallel !$acc parallel self (.true.) @@ -36,6 +40,8 @@ program test !$acc parallel self (i.gt.1) !$acc end parallel + !$acc kernels self + !$acc end kernels !$acc kernels self (x) !$acc end kernels !$acc kernels self (.true.) @@ -43,6 +49,8 @@ program test !$acc kernels self (i.gt.1) !$acc end kernels + !$acc serial self + !$acc end serial !$acc serial self (x) !$acc end serial !$acc serial self (.true.) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c index 5398905f411..f04c02fdb89 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c @@ -1,3 +1,7 @@ +/* OpenACC 'if' clause. */ + +/* See also 'self-1.c'. */ + #include #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c new file mode 100644 index 00000000000..f77edda83e3 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c @@ -0,0 +1,36 @@ +/* Test 'if' and 'self' clause appearing together. */ + +#include + +static int test(float i, long double s) +{ + int ret; +#pragma acc serial copyout(ret) if(i) self(s) + /* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { xfail openacc_nvidia_accel_selected } .-1 } */ + { + ret = acc_on_device(acc_device_host); + } + return ret; +} + +int main() +{ + if (!test(0, 0)) + __builtin_abort(); + + if (!test(0, 1)) + __builtin_abort(); + +#if ACC_MEM_SHARED + if (!test(1, 0)) + __builtin_abort(); +#else + if (test(1, 0)) + __builtin_abort(); +#endif + + if (!test(1, 1)) + __builtin_abort(); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c index 752e16e8545..d17e27c8af3 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c @@ -1,3 +1,8 @@ +/* OpenACC 'self' clause. */ + +/* This is 'if-1.c' with 'self(!cond)' instead of 'if(cond)' on compute + constructs. */ + #include #include #include diff --git a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 index e0cfd912d0f..5ba6509749e 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 @@ -1,3 +1,7 @@ +! OpenACC 'if' clause. + +! See also 'self-1.f90'. + ! { dg-do run } ! { dg-additional-options "-cpp" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 new file mode 100644 index 00000000000..b9ec9de08d9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 @@ -0,0 +1,996 @@ +! OpenACC 'self' clause. + +! This is 'if-1.f90' with 'self(!cond)' instead of 'if(cond)' on compute +! constructs. +! ..., which the exception of certain 'kernels' constructs. + +! { dg-do run } +! { dg-additional-options "-cpp" } + +! { dg-additional-options "--param=openacc-kernels=decompose" } + +! { dg-additional-options "-fopt-info-note-omp" } +! { dg-additional-options "-foffload=-fopt-info-note-omp" } + +! { dg-additional-options "--param=openacc-privatization=noisy" } +! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types): +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } + +! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' +! passed to 'incr' may be unset, and in that case, it will be set to [...]", +! so to maintain compatibility with earlier Tcl releases, we manually +! initialize counter variables: +! { dg-line l_dummy[variable c_compute 0] } +! { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid +! "WARNING: dg-line var l_dummy defined, but not used". */ + +program main + use openacc + implicit none + + integer, parameter :: N = 8 + integer, parameter :: one = 1 + integer, parameter :: zero = 0 + integer i, nn + real, allocatable :: a(:), b(:) + real exp, exp2 + + i = 0 + + allocate (a(N)) + allocate (b(N)) + + a(:) = 4.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (1 /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + !TODO Unhandled 'CONST_DECL' instances for constant argument in 'acc_on_device' call. + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 5.0 +#else + exp = 4.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 1 + end do + + a(:) = 16.0 + + !$acc parallel self (0 /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 17.0) STOP 2 + end do + + a(:) = 8.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 9.0 +#else + exp = 8.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 3 + end do + + a(:) = 22.0 + + !$acc parallel self (zero /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 23.0) STOP 4 + end do + + a(:) = 16.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.FALSE.) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 17.0; +#else + exp = 16.0; +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 5 + end do + + a(:) = 76.0 + + !$acc parallel self (.TRUE.) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 77.0) STOP 6 + end do + + a(:) = 22.0 + + nn = 1 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (nn /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 23.0; +#else + exp = 22.0; +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 7 + end do + + a(:) = 18.0 + + nn = 0 + + !$acc parallel self (nn /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 19.0) STOP 8 + end do + + a(:) = 49.0 + + nn = 1 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.NOT. ((nn + nn) > 0)) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 50.0 +#else + exp = 49.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 9 + end do + + a(:) = 38.0 + + nn = 0; + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.NOT. ((nn + nn) > 0)) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 39.0) STOP 10 + end do + + a(:) = 91.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.NOT. (-2 > 0)) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 92.0) STOP 11 + end do + + a(:) = 43.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 44.0 +#else + exp = 43.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 12 + end do + + a(:) = 87.0 + + !$acc parallel self (one /= 0) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 88.0) STOP 13 + end do + + a(:) = 3.0 + b(:) = 9.0 + +#if ACC_MEM_SHARED + exp = 0.0 + exp2 = 0.0 +#else + call acc_copyin (a, sizeof (a)) + call acc_copyin (b, sizeof (b)) + exp = 3.0; + exp2 = 9.0; +#endif + + !$acc update device (a(1:N), b(1:N)) if (1 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (1 == 1) + + do i = 1, N + if (a(i) .ne. exp) STOP 14 + if (b(i) .ne. exp2) STOP 15 + end do + + a(:) = 6.0 + b(:) = 12.0 + + !$acc update device (a(1:N), b(1:N)) if (0 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (1 == 1) + + do i = 1, N + if (a(i) .ne. exp) STOP 16 + if (b(i) .ne. exp2) STOP 17 + end do + + a(:) = 26.0 + b(:) = 21.0 + + !$acc update device (a(1:N), b(1:N)) if (1 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (0 == 1) + + do i = 1, N + if (a(i) .ne. 0.0) STOP 18 + if (b(i) .ne. 0.0) STOP 19 + end do + +#if !ACC_MEM_SHARED + call acc_copyout (a, sizeof (a)) + call acc_copyout (b, sizeof (b)) +#endif + + a(:) = 4.0 + b(:) = 0.0 + + !$acc data copyin (a(1:N)) copyout (b(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + + !$acc parallel present (a(1:N)) + do i = 1, N + b(i) = a(i) + end do + !$acc end parallel + !$acc end data + + do i = 1, N + if (b(i) .ne. 4.0) STOP 20 + end do + + a(:) = 8.0 + b(:) = 1.0 + + !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 } + +#if !ACC_MEM_SHARED + if (acc_is_present (a) .eqv. .TRUE.) STOP 21 + if (acc_is_present (b) .eqv. .TRUE.) STOP 22 +#endif + + !$acc end data + + a(:) = 18.0 + b(:) = 21.0 + + !$acc data copyin (a(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } + +#if !ACC_MEM_SHARED + if (acc_is_present (a) .eqv. .FALSE.) STOP 23 +#endif + + !$acc data copyout (b(1:N)) if (0 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 24 +#endif + !$acc data copyout (b(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + + !$acc parallel present (a(1:N)) present (b(1:N)) + do i = 1, N + b(i) = a(i) + end do + !$acc end parallel + + !$acc end data + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 25 +#endif + !$acc end data + !$acc end data + + do i = 1, N + if (b(1) .ne. 18.0) STOP 26 + end do + + !$acc enter data copyin (b(1:N)) if (0 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 27 +#endif + + !$acc exit data delete (b(1:N)) if (0 == 1) + + !$acc enter data copyin (b(1:N)) if (1 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 28 +#endif + + !$acc exit data delete (b(1:N)) if (1 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 29 +#endif + + !$acc enter data copyin (b(1:N)) if (zero == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 30 +#endif + + !$acc exit data delete (b(1:N)) if (zero == 1) + + !$acc enter data copyin (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 31 +#endif + + !$acc exit data delete (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 32 +#endif + + !$acc enter data copyin (b(1:N)) if (one == 0) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 33 +#endif + + !$acc exit data delete (b(1:N)) if (one == 0) + + !$acc enter data copyin (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 34 +#endif + + !$acc exit data delete (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 35 +#endif + + a(:) = 4.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (1 /= 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 5.0 +#else + exp = 4.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 36 + end do + + a(:) = 16.0 + + !$acc kernels if (0 == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 17.0) STOP 37 + end do + + a(:) = 8.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 9.0 +#else + exp = 8.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 38 + end do + + a(:) = 22.0 + + !$acc kernels if (zero == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 23.0) STOP 39 + end do + + a(:) = 16.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (.FALSE.) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 17.0; +#else + exp = 16.0; +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 40 + end do + + a(:) = 76.0 + + !$acc kernels if (.FALSE.) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 77.0) STOP 41 + end do + + a(:) = 22.0 + + nn = 1 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (nn /= 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 23.0; +#else + exp = 22.0; +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 42 + end do + + a(:) = 18.0 + + nn = 0 + + !$acc kernels if (nn == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 19.0) STOP 43 + end do + + a(:) = 49.0 + + nn = 1 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (.NOT. ((nn + nn) > 0)) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 50.0 +#else + exp = 49.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 44 + end do + + a(:) = 38.0 + + nn = 0; + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if ((nn + nn) > 0) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 39.0) STOP 45 + end do + + a(:) = 91.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (-2 > 0) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 92.0) STOP 46 + end do + + a(:) = 43.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 44.0 +#else + exp = 43.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 47 + end do + + a(:) = 87.0 + + !$acc kernels if (one == 0) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 88.0) STOP 48 + end do + + a(:) = 3.0 + b(:) = 9.0 + +#if ACC_MEM_SHARED + exp = 0.0 + exp2 = 0.0 +#else + call acc_copyin (a, sizeof (a)) + call acc_copyin (b, sizeof (b)) + exp = 3.0; + exp2 = 9.0; +#endif + + !$acc update device (a(1:N), b(1:N)) if (1 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (1 == 1) + + do i = 1, N + if (a(i) .ne. exp) STOP 49 + if (b(i) .ne. exp2) STOP 50 + end do + + a(:) = 6.0 + b(:) = 12.0 + + !$acc update device (a(1:N), b(1:N)) if (0 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (1 == 1) + + do i = 1, N + if (a(i) .ne. exp) STOP 51 + if (b(i) .ne. exp2) STOP 52 + end do + + a(:) = 26.0 + b(:) = 21.0 + + !$acc update device (a(1:N), b(1:N)) if (1 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (0 == 1) + + do i = 1, N + if (a(i) .ne. 0.0) STOP 53 + if (b(i) .ne. 0.0) STOP 54 + end do + +#if !ACC_MEM_SHARED + call acc_copyout (a, sizeof (a)) + call acc_copyout (b, sizeof (b)) +#endif + + a(:) = 4.0 + b(:) = 0.0 + + !$acc data copyin (a(1:N)) copyout (b(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + + !$acc kernels present (a(1:N)) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + b(i) = a(i) + end do + !$acc end kernels + !$acc end data + + do i = 1, N + if (b(i) .ne. 4.0) STOP 55 + end do + + a(:) = 8.0 + b(:) = 1.0 + + !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 } + +#if !ACC_MEM_SHARED + if (acc_is_present (a) .eqv. .TRUE.) STOP 56 + if (acc_is_present (b) .eqv. .TRUE.) STOP 57 +#endif + + !$acc end data + + a(:) = 18.0 + b(:) = 21.0 + + !$acc data copyin (a(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } + +#if !ACC_MEM_SHARED + if (acc_is_present (a) .eqv. .FALSE.) STOP 58 +#endif + + !$acc data copyout (b(1:N)) if (0 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 59 +#endif + !$acc data copyout (b(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + + !$acc kernels present (a(1:N)) present (b(1:N)) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + b(i) = a(i) + end do + !$acc end kernels + + !$acc end data + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 60 +#endif + !$acc end data + !$acc end data + + do i = 1, N + if (b(1) .ne. 18.0) STOP 61 + end do + + !$acc enter data copyin (b(1:N)) if (0 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 62 +#endif + + !$acc exit data delete (b(1:N)) if (0 == 1) + + !$acc enter data copyin (b(1:N)) if (1 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 63 +#endif + + !$acc exit data delete (b(1:N)) if (1 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 64 +#endif + + !$acc enter data copyin (b(1:N)) if (zero == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 65 +#endif + + !$acc exit data delete (b(1:N)) if (zero == 1) + + !$acc enter data copyin (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 66 +#endif + + !$acc exit data delete (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 67 +#endif + + !$acc enter data copyin (b(1:N)) if (one == 0) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 68 +#endif + + !$acc exit data delete (b(1:N)) if (one == 0) + + !$acc enter data copyin (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 69 +#endif + + !$acc exit data delete (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 70 +#endif + +end program main -- 2.34.1