From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: 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 Content-Type: multipart/mixed; boundary="=-=-=" 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 List-Id: --=-=-= Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable 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 construct= s > .../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 construct= s", 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=C3=BC=C3=9Fe Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstra=C3=9Fe 201= , 80634 M=C3=BCnchen; Gesellschaft mit beschr=C3=A4nkter Haftung; Gesch=C3= =A4ftsf=C3=BChrer: Thomas Heurung, Frank Th=C3=BCrauf; Sitz der Gesellschaf= t: M=C3=BCnchen; Registergericht M=C3=BCnchen, HRB 106955 --=-=-= Content-Type: text/x-diff Content-Disposition: inline; filename="0001-Extend-test-suite-coverage-for-OpenACC-self-clause-f.patch" >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 --=-=-=--