From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 72886 invoked by alias); 13 Dec 2018 14:14:24 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 72817 invoked by uid 89); 13 Dec 2018 14:14:17 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-26.9 required=5.0 tests=BAYES_00,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,RCVD_IN_DNSWL_NONE,SPF_PASS,TIME_LIMIT_EXCEEDED autolearn=unavailable version=3.3.2 spammy=1111, UD:b, gang, 15,5 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Thu, 13 Dec 2018 14:14:06 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-MBX-04.mgc.mentorg.com) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1gXRka-00017B-EK from Julian_Brown@mentor.com ; Thu, 13 Dec 2018 06:14:04 -0800 Received: from squid.athome (137.202.0.90) by SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Thu, 13 Dec 2018 14:14:00 +0000 Date: Thu, 13 Dec 2018 14:14:00 -0000 From: Julian Brown To: Jakub Jelinek CC: Cesar Philippidis , Thomas Schwinge , "gcc-patches@gcc.gnu.org" , Tom de Vries Subject: Re: [patch] various OpenACC reduction enhancements - test cases Message-ID: <20181213141358.72e4419f@squid.athome> In-Reply-To: <20181204125933.GM12380@tucnak> References: <20181204125933.GM12380@tucnak> MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="MP_/.H_X2eEN/RJX2_pabxIevKR" X-IsSubscribed: yes X-SW-Source: 2018-12/txt/msg00931.txt.bz2 --MP_/.H_X2eEN/RJX2_pabxIevKR Content-Type: text/plain; charset="US-ASCII" Content-Transfer-Encoding: 7bit Content-Disposition: inline Content-length: 1243 On Tue, 4 Dec 2018 13:59:33 +0100 Jakub Jelinek wrote: > On Fri, Jun 29, 2018 at 11:23:21AM -0700, Cesar Philippidis wrote: > > Attached are the updated reductions tests cases. Again, these have > > been bootstrapped and regression tested cleanly for x86_64 with > > nvptx offloading. Is it OK for trunk? > > If Thomas is ok with this, it is ok for trunk. Here's a new version to go with the FE patch posted here: https://gcc.gnu.org/ml/gcc-patches/2018-12/msg00930.html Thanks, Julian ChangeLog 2018-xx-xx Cesar Philippidis Nathan Sidwell Julian Brown gcc/testsuite/ * c-c++-common/goacc/orphan-reductions-1.c: New test. * c-c++-common/goacc/reduction-7.c: New test. * c-c++-common/goacc/routine-4.c: Update. * g++.dg/goacc/reductions-1.C: New test. * gcc.dg/goacc/loop-processing-1.c: Update. * gfortran.dg/goacc/orphan-reductions-1.f90: New test. libgomp/ * libgomp.oacc-c-c++-common/par-reduction-3.c: New test. * libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c: New test. * libgomp.oacc-fortran/reduction-9.f90: New test. --MP_/.H_X2eEN/RJX2_pabxIevKR Content-Type: text/x-patch Content-Transfer-Encoding: quoted-printable Content-Disposition: attachment; filename="trunk-reductions-tests-2.diff" Content-length: 35402 commit 7d445a56d6db96696cec8359e58258d47fa7c9ae Author: Julian Brown Date: Wed Dec 12 11:11:03 2018 -0800 Various OpenACC reduction enhancements - test cases =20=20=20=20 2018-xx-xx Cesar Philippidis Nathan Sidwell Julian Brown =20=20=20=20 gcc/testsuite/ * c-c++-common/goacc/orphan-reductions-1.c: New test. * c-c++-common/goacc/reduction-7.c: New test. * c-c++-common/goacc/routine-4.c: Update. * g++.dg/goacc/reductions-1.C: New test. * gcc.dg/goacc/loop-processing-1.c: Update. * gfortran.dg/goacc/orphan-reductions-1.f90: New test. =20=20=20=20 libgomp/ * libgomp.oacc-c-c++-common/par-reduction-3.c: New test. * libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c: New test. * libgomp.oacc-fortran/reduction-9.f90: New test. diff --git a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-1.c b/gcc/t= estsuite/c-c++-common/goacc/orphan-reductions-1.c new file mode 100644 index 0000000..b0bd4a7 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-1.c @@ -0,0 +1,56 @@ +/* Test orphan reductions. */ + +#include + +#pragma acc routine seq +int +seq_reduction (int n) +{ + int i, sum =3D 0; +#pragma acc loop seq reduction(+:sum) + for (i =3D 0; i < n; i++) + sum =3D sum + 1; + + return sum; +} + +#pragma acc routine gang +int +gang_reduction (int n) +{ + int i, s1 =3D 0, s2 =3D 0; +#pragma acc loop gang reduction(+:s1) /* { dg-error "gang reduction on an = orphan loop" } */ + for (i =3D 0; i < n; i++) + s1 =3D s1 + 2; + +#pragma acc loop gang reduction(+:s2) /* { dg-error "gang reduction on an = orphan loop" } */ + for (i =3D 0; i < n; i++) + s2 =3D s2 + 2; + + + return s1 + s2; +} + +#pragma acc routine worker +int +worker_reduction (int n) +{ + int i, sum =3D 0; +#pragma acc loop worker reduction(+:sum) + for (i =3D 0; i < n; i++) + sum =3D sum + 3; + + return sum; +} + +#pragma acc routine vector +int +vector_reduction (int n) +{ + int i, sum =3D 0; +#pragma acc loop vector reduction(+:sum) + for (i =3D 0; i < n; i++) + sum =3D sum + 4; + + return sum; +} diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-7.c b/gcc/testsuite= /c-c++-common/goacc/reduction-7.c new file mode 100644 index 0000000..eba1d02 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/reduction-7.c @@ -0,0 +1,111 @@ +/* Exercise invalid reductions on array and struct members. */ + +void +test_parallel () +{ + struct { + int a; + float b[5]; + } s1, s2[10]; + + int i; + double z[100]; + +#pragma acc parallel reduction(+:s1.a) /* { dg-error "expected '\\\)' befo= re '\\\.' token" } */ + for (i =3D 0; i < 10; i++) + s1.a +=3D 1; + +#pragma acc parallel reduction(+:s1.b[3]) /* { dg-error "expected '\\\)' b= efore '\\\.' token" } */ + for (i =3D 0; i < 10; i++) + s1.b[3] +=3D 1; + +#pragma acc parallel reduction(+:s2[2].a) /* { dg-error "expected '\\\)' b= efore '\\\[' token" } */ + for (i =3D 0; i < 10; i++) + s2[2].a +=3D 1; + +#pragma acc parallel reduction(+:s2[3].b[4]) /* { dg-error "expected '\\\)= ' before '\\\[' token" } */ + for (i =3D 0; i < 10; i++) + s2[3].b[4] +=3D 1; + +#pragma acc parallel reduction(+:z[5]) /* { dg-error "expected '\\\)' befo= re '\\\[' token" } */ + for (i =3D 0; i < 10; i++) + z[5] +=3D 1; +} + +void +test_combined () +{ + struct { + int a; + float b[5]; + } s1, s2[10]; + + int i; + double z[100]; + +#pragma acc parallel loop reduction(+:s1.a) /* { dg-error "expected '\\\)'= before '\\\.' token" } */ + for (i =3D 0; i < 10; i++) + s1.a +=3D 1; + +#pragma acc parallel loop reduction(+:s1.b[3]) /* { dg-error "expected '\\= \)' before '\\\.' token" } */ + for (i =3D 0; i < 10; i++) + s1.b[3] +=3D 1; + +#pragma acc parallel loop reduction(+:s2[2].a) /* { dg-error "expected '\\= \)' before '\\\[' token" } */ + for (i =3D 0; i < 10; i++) + s2[2].a +=3D 1; + +#pragma acc parallel loop reduction(+:s2[3].b[4]) /* { dg-error "expected = '\\\)' before '\\\[' token" } */ + for (i =3D 0; i < 10; i++) + s2[3].b[4] +=3D 1; + +#pragma acc parallel loop reduction(+:z[5]) /* { dg-error "expected '\\\)'= before '\\\[' token" } */ + for (i =3D 0; i < 10; i++) + z[5] +=3D 1; + +} + +void +test_loops () +{ + struct { + int a; + float b[5]; + } s1, s2[10]; + + int i; + double z[100]; + +#pragma acc parallel + { +#pragma acc loop reduction(+:s1.a) /* { dg-error "expected '\\\)' before '= \\\.' token" } */ + for (i =3D 0; i < 10; i++) + s1.a +=3D 1; + +#pragma acc loop reduction(+:s1.b[3]) /* { dg-error "expected '\\\)' befor= e '\\\.' token" } */ + for (i =3D 0; i < 10; i++) + s1.b[3] +=3D 1; + +#pragma acc loop reduction(+:s2[2].a) /* { dg-error "expected '\\\)' befor= e '\\\[' token" } */ + for (i =3D 0; i < 10; i++) + s2[2].a +=3D 1; + +#pragma acc loop reduction(+:s2[3].b[4]) /* { dg-error "expected '\\\)' be= fore '\\\[' token" } */ + for (i =3D 0; i < 10; i++) + s2[3].b[4] +=3D 1; + +#pragma acc loop reduction(+:z[5]) /* { dg-error "expected '\\\)' before '= \\\[' token" } */ + for (i =3D 0; i < 10; i++) + z[5] +=3D 1; + } +} + +int +main () +{ + test_parallel (); + test_combined (); + test_loops (); + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/routine-4.c b/gcc/testsuite/c= -c++-common/goacc/routine-4.c index efc4a0b..91abfb5 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-4.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-4.c @@ -22,7 +22,7 @@ void seq (void) for (int i =3D 0; i < 10; i++) red ++; =20 -#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by conta= ining routine" } +#pragma acc loop seq reduction (+:red) for (int i =3D 0; i < 10; i++) red ++; =20 @@ -48,7 +48,7 @@ void vector (void) /* { dg-message "declared here" "1" } = */ for (int i =3D 0; i < 10; i++) red ++; =20 -#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by conta= ining routine" } +#pragma acc loop seq reduction (+:red) for (int i =3D 0; i < 10; i++) red ++; =20 @@ -74,7 +74,7 @@ void worker (void) /* { dg-message "declared here" "2" } = */ for (int i =3D 0; i < 10; i++) red ++; =20 -#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by conta= ining routine" } +#pragma acc loop seq reduction (+:red) for (int i =3D 0; i < 10; i++) red ++; =20 @@ -100,7 +100,7 @@ void gang (void) /* { dg-message "declared here" "3" } = */ for (int i =3D 0; i < 10; i++) red ++; =20 -#pragma acc loop gang reduction (+:red) +#pragma acc loop seq reduction (+:red) for (int i =3D 0; i < 10; i++) red ++; =20 diff --git a/gcc/testsuite/g++.dg/goacc/reductions-1.C b/gcc/testsuite/g++.= dg/goacc/reductions-1.C new file mode 100644 index 0000000..18f43f4 --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/reductions-1.C @@ -0,0 +1,548 @@ +// Test for invalid reduction variables. + +class C1 +{ + int b, d[10]; + +public: + int a, c[10]; + + C1 () { a =3D 0; b =3D 0; } + int& get_b () { return b; } + int* get_d () { return d; } +}; + +template +class C2 +{ + T b, d[10]; + +public: + T a, c[10]; + + C2 () { a =3D 0; b =3D 0; } + T& get_b () { return b; } + T* get_d () { return d; } +}; + +struct S1 +{ + int a, b, c[10], d[10]; + + S1 () { a =3D 0; b =3D 0; } + int& get_b () { return b; } + int* get_d () { return d; } +}; + +template +struct S2 +{ + T a, b, c[10], d[10]; + + S2 () { a =3D 0; b =3D 0; } + T& get_b () { return b; } + T* get_d () { return d; } +}; + +template +void +test_parallel () +{ + int i, a[10]; + T b[10]; + C1 c1, c1a[10]; + C2 c2, c2a[10]; + S1 s1, s1a[10]; + S2 s2, s2a[10]; + + // Reductions on class members. + +#pragma acc parallel reduction(+:c1.a) // { dg-error "expected '\\\)' befo= re '\\\.' token" } + for (i =3D 0; i < 100; i++) + c1.a +=3D 1; + +#pragma acc parallel reduction(+:c1.get_b ()) // { dg-error "expected '\\\= )' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + c1.get_b () +=3D 1; + +#pragma acc parallel reduction(+:c1.c[1]) // { dg-error "expected '\\\)' b= efore '\\\.' token" } + for (i =3D 0; i < 100; i++) + c1.c[1] +=3D 1; + +#pragma acc parallel reduction(+:c1.get_d ()[1]) // { dg-error "expected '= \\\)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + c1.get_d ()[1] +=3D 1; + +#pragma acc parallel reduction(+:c1a[1].a) // { dg-error "expected '\\\)' = before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c1a[1].a +=3D 1; + +#pragma acc parallel reduction(+:c1a[1].get_b ()) // { dg-error "expected = '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c1a[1].get_b () +=3D 1; + +#pragma acc parallel reduction(+:c1a[1].c[1]) // { dg-error "expected '\\\= )' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c1a[1].c[1] +=3D 1; + +#pragma acc parallel reduction(+:c1a[1].get_d ()[1]) // { dg-error "expect= ed '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c1a[1].get_d ()[1] +=3D 1; + + + // Reductions on a template class member. + +#pragma acc parallel reduction(+:c2.a) // { dg-error "expected '\\\)' befo= re '\\\.' token" } + for (i =3D 0; i < 100; i++) + c2.a +=3D 1; + +#pragma acc parallel reduction(+:c2.get_b ()) // { dg-error "expected '\\\= )' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + c2.get_b () +=3D 1; + +#pragma acc parallel reduction(+:c2.c[1]) // { dg-error "expected '\\\)' b= efore '\\\.' token" } + for (i =3D 0; i < 100; i++) + c2.c[1] +=3D 1; + +#pragma acc parallel reduction(+:c2.get_d ()[1]) // { dg-error "expected '= \\\)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + c2.get_d ()[1] +=3D 1; + + +#pragma acc parallel reduction(+:c2a[1].a) // { dg-error "expected '\\\)' = before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c2a[1].a +=3D 1; + +#pragma acc parallel reduction(+:c2a[1].get_b ()[1]) // { dg-error "expect= ed '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c2a[1].get_b () +=3D 1; + +#pragma acc parallel reduction(+:c2a[1].c[1]) // { dg-error "expected '\\\= )' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c2a[1].c[1] +=3D 1; + +#pragma acc parallel reduction(+:c2a[1].get_d ()[1]) // { dg-error "expect= ed '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c2a[1].get_d ()[1] +=3D 1; + + + // Reductions on struct element. + +#pragma acc parallel reduction(+:s1.a) // { dg-error "expected '\\\)' befo= re '\\\.' token" } + for (i =3D 0; i < 100; i++) + s1.a +=3D 1; + +#pragma acc parallel reduction(+:s1.get_b ()) // { dg-error "expected '\\\= )' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + s1.get_b () +=3D 1; + +#pragma acc parallel reduction(+:s1.c[1]) // { dg-error "expected '\\\)' b= efore '\\\.' token" } + for (i =3D 0; i < 100; i++) + s1.c[1] +=3D 1; + +#pragma acc parallel reduction(+:s1.get_d ()[1]) // { dg-error "expected '= \\\)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + s1.get_d ()[1] +=3D 1; + +#pragma acc parallel reduction(+:s1a[1].a) // { dg-error "expected '\\\)' = before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s1a[1].a +=3D 1; + +#pragma acc parallel reduction(+:s1a[1].get_b ()) // { dg-error "expected = '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s1a[1].get_b () +=3D 1; + +#pragma acc parallel reduction(+:s1a[1].c[1]) // { dg-error "expected '\\\= )' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s1a[1].c[1] +=3D 1; + +#pragma acc parallel reduction(+:s1a[1].get_d ()[1]) // { dg-error "expect= ed '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s1a[1].get_d ()[1] +=3D 1; + + + // Reductions on a template struct element. + +#pragma acc parallel reduction(+:s2.a) // { dg-error "expected '\\\)' befo= re '\\\.' token" } + for (i =3D 0; i < 100; i++) + s2.a +=3D 1; + +#pragma acc parallel reduction(+:s2.get_b ()) // { dg-error "expected '\\\= )' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + s2.get_b () +=3D 1; + +#pragma acc parallel reduction(+:s2.c[1]) // { dg-error "expected '\\\)' b= efore '\\\.' token" } + for (i =3D 0; i < 100; i++) + s2.c[1] +=3D 1; + +#pragma acc parallel reduction(+:s2.get_d ()[1]) // { dg-error "expected '= \\\)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + s2.get_d ()[1] +=3D 1; + +#pragma acc parallel reduction(+:s2a[1].a) // { dg-error "expected '\\\)' = before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s2a[1].a +=3D 1; + +#pragma acc parallel reduction(+:s2a[1].get_b ()) // { dg-error "expected = '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s2a[1].get_b () +=3D 1; + +#pragma acc parallel reduction(+:s2a[1].c[1]) // { dg-error "expected '\\\= )' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s2a[1].c[1] +=3D 1; + +#pragma acc parallel reduction(+:s2a[1].get_d ()[1]) // { dg-error "expect= ed '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s2a[1].get_d ()[1] +=3D 1; + + + // Reductions on arrays. + +#pragma acc parallel reduction(+:a[10]) // { dg-error "expected '\\\)' bef= ore '\\\[' token" } + for (i =3D 0; i < 100; i++) + a[10] +=3D 1; + +#pragma acc parallel reduction(+:b[10]) // { dg-error "expected '\\\)' bef= ore '\\\[' token" } + for (i =3D 0; i < 100; i++) + b[10] +=3D 1; +} + +template +void +test_combined () +{ + int i, a[10]; + T b[10]; + C1 c1, c1a[10]; + C2 c2, c2a[10]; + S1 s1, s1a[10]; + S2 s2, s2a[10]; + + // Reductions on class members. + +#pragma acc parallel loop reduction(+:c1.a) // { dg-error "expected '\\\)'= before '\\\.' token" } + for (i =3D 0; i < 100; i++) + c1.a +=3D 1; + +#pragma acc parallel loop reduction(+:c1.get_b ()) // { dg-error "expected= '\\\)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + c1.get_b () +=3D 1; + +#pragma acc parallel loop reduction(+:c1.c[1]) // { dg-error "expected '\\= \)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + c1.c[1] +=3D 1; + +#pragma acc parallel loop reduction(+:c1.get_d ()[1]) // { dg-error "expec= ted '\\\)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + c1.get_d ()[1] +=3D 1; + +#pragma acc parallel loop reduction(+:c1a[1].a) // { dg-error "expected '\= \\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c1a[1].a +=3D 1; + +#pragma acc parallel loop reduction(+:c1a[1].get_b ()) // { dg-error "expe= cted '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c1a[1].get_b () +=3D 1; + +#pragma acc parallel loop reduction(+:c1a[1].c[1]) // { dg-error "expected= '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c1a[1].c[1] +=3D 1; + +#pragma acc parallel loop reduction(+:c1a[1].get_d ()[1]) // { dg-error "e= xpected '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c1a[1].get_d ()[1] +=3D 1; + + + // Reductions on a template class member. + +#pragma acc parallel loop reduction(+:c2.a) // { dg-error "expected '\\\)'= before '\\\.' token" } + for (i =3D 0; i < 100; i++) + c2.a +=3D 1; + +#pragma acc parallel loop reduction(+:c2.get_b ()) // { dg-error "expected= '\\\)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + c2.get_b () +=3D 1; + +#pragma acc parallel loop reduction(+:c2.c[1]) // { dg-error "expected '\\= \)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + c2.c[1] +=3D 1; + +#pragma acc parallel loop reduction(+:c2.get_d ()[1]) // { dg-error "expec= ted '\\\)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + c2.get_d ()[1] +=3D 1; + + +#pragma acc parallel loop reduction(+:c2a[1].a) // { dg-error "expected '\= \\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c2a[1].a +=3D 1; + +#pragma acc parallel loop reduction(+:c2a[1].get_b ()[1]) // { dg-error "e= xpected '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c2a[1].get_b () +=3D 1; + +#pragma acc parallel loop reduction(+:c2a[1].c[1]) // { dg-error "expected= '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c2a[1].c[1] +=3D 1; + +#pragma acc parallel loop reduction(+:c2a[1].get_d ()[1]) // { dg-error "e= xpected '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c2a[1].get_d ()[1] +=3D 1; + + + // Reductions on struct element. + +#pragma acc parallel loop reduction(+:s1.a) // { dg-error "expected '\\\)'= before '\\\.' token" } + for (i =3D 0; i < 100; i++) + s1.a +=3D 1; + +#pragma acc parallel loop reduction(+:s1.get_b ()) // { dg-error "expected= '\\\)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + s1.get_b () +=3D 1; + +#pragma acc parallel loop reduction(+:s1.c[1]) // { dg-error "expected '\\= \)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + s1.c[1] +=3D 1; + +#pragma acc parallel loop reduction(+:s1.get_d ()[1]) // { dg-error "expec= ted '\\\)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + s1.get_d ()[1] +=3D 1; + +#pragma acc parallel loop reduction(+:s1a[1].a) // { dg-error "expected '\= \\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s1a[1].a +=3D 1; + +#pragma acc parallel loop reduction(+:s1a[1].get_b ()) // { dg-error "expe= cted '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s1a[1].get_b () +=3D 1; + +#pragma acc parallel loop reduction(+:s1a[1].c[1]) // { dg-error "expected= '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s1a[1].c[1] +=3D 1; + +#pragma acc parallel loop reduction(+:s1a[1].get_d ()[1]) // { dg-error "e= xpected '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s1a[1].get_d ()[1] +=3D 1; + + + // Reductions on a template struct element. + +#pragma acc parallel loop reduction(+:s2.a) // { dg-error "expected '\\\)'= before '\\\.' token" } + for (i =3D 0; i < 100; i++) + s2.a +=3D 1; + +#pragma acc parallel loop reduction(+:s2.get_b ()) // { dg-error "expected= '\\\)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + s2.get_b () +=3D 1; + +#pragma acc parallel loop reduction(+:s2.c[1]) // { dg-error "expected '\\= \)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + s2.c[1] +=3D 1; + +#pragma acc parallel loop reduction(+:s2.get_d ()[1]) // { dg-error "expec= ted '\\\)' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + s2.get_d ()[1] +=3D 1; + +#pragma acc parallel loop reduction(+:s2a[1].a) // { dg-error "expected '\= \\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s2a[1].a +=3D 1; + +#pragma acc parallel loop reduction(+:s2a[1].get_b ()) // { dg-error "expe= cted '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s2a[1].get_b () +=3D 1; + +#pragma acc parallel loop reduction(+:s2a[1].c[1]) // { dg-error "expected= '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s2a[1].c[1] +=3D 1; + +#pragma acc parallel loop reduction(+:s2a[1].get_d ()[1]) // { dg-error "e= xpected '\\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s2a[1].get_d ()[1] +=3D 1; + + + // Reductions on arrays. + +#pragma acc parallel loop reduction(+:a[10]) // { dg-error "expected '\\\)= ' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + a[10] +=3D 1; + +#pragma acc parallel loop reduction(+:b[10]) // { dg-error "expected '\\\)= ' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + b[10] +=3D 1; +} + +template +void +test_loop () +{ + int i, a[10]; + T b[10]; + C1 c1, c1a[10]; + C2 c2, c2a[10]; + S1 s1, s1a[10]; + S2 s2, s2a[10]; + + // Reductions on class members. + + #pragma acc parallel + { + +#pragma acc loop reduction(+:c1.a) // { dg-error "expected '\\\)' before '= \\\.' token" } + for (i =3D 0; i < 100; i++) + c1.a +=3D 1; + +#pragma acc loop reduction(+:c1.get_b ()) // { dg-error "expected '\\\)' b= efore '\\\.' token" } + for (i =3D 0; i < 100; i++) + c1.get_b () +=3D 1; + +#pragma acc loop reduction(+:c1.c[1]) // { dg-error "expected '\\\)' befor= e '\\\.' token" } + for (i =3D 0; i < 100; i++) + c1.c[1] +=3D 1; + +#pragma acc loop reduction(+:c1.get_d ()[1]) // { dg-error "expected '\\\)= ' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + c1.get_d ()[1] +=3D 1; + +#pragma acc loop reduction(+:c1a[1].a) // { dg-error "expected '\\\)' befo= re '\\\[' token" } + for (i =3D 0; i < 100; i++) + c1a[1].a +=3D 1; + +#pragma acc loop reduction(+:c1a[1].get_b ()) // { dg-error "expected '\\\= )' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c1a[1].get_b () +=3D 1; + +#pragma acc loop reduction(+:c1a[1].c[1]) // { dg-error "expected '\\\)' b= efore '\\\[' token" } + for (i =3D 0; i < 100; i++) + c1a[1].c[1] +=3D 1; + +#pragma acc loop reduction(+:c1a[1].get_d ()[1]) // { dg-error "expected '= \\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c1a[1].get_d ()[1] +=3D 1; + + + // Reductions on a template class member. + +#pragma acc loop reduction(+:c2.a) // { dg-error "expected '\\\)' before '= \\\.' token" } + for (i =3D 0; i < 100; i++) + c2.a +=3D 1; + +#pragma acc loop reduction(+:c2.get_b ()) // { dg-error "expected '\\\)' b= efore '\\\.' token" } + for (i =3D 0; i < 100; i++) + c2.get_b () +=3D 1; + +#pragma acc loop reduction(+:c2.c[1]) // { dg-error "expected '\\\)' befor= e '\\\.' token" } + for (i =3D 0; i < 100; i++) + c2.c[1] +=3D 1; + +#pragma acc loop reduction(+:c2.get_d ()[1]) // { dg-error "expected '\\\)= ' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + c2.get_d ()[1] +=3D 1; + + +#pragma acc loop reduction(+:c2a[1].a) // { dg-error "expected '\\\)' befo= re '\\\[' token" } + for (i =3D 0; i < 100; i++) + c2a[1].a +=3D 1; + +#pragma acc loop reduction(+:c2a[1].get_b ()[1]) // { dg-error "expected '= \\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c2a[1].get_b () +=3D 1; + +#pragma acc loop reduction(+:c2a[1].c[1]) // { dg-error "expected '\\\)' b= efore '\\\[' token" } + for (i =3D 0; i < 100; i++) + c2a[1].c[1] +=3D 1; + +#pragma acc loop reduction(+:c2a[1].get_d ()[1]) // { dg-error "expected '= \\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + c2a[1].get_d ()[1] +=3D 1; + + + // Reductions on struct element. + +#pragma acc loop reduction(+:s1.a) // { dg-error "expected '\\\)' before '= \\\.' token" } + for (i =3D 0; i < 100; i++) + s1.a +=3D 1; + +#pragma acc loop reduction(+:s1.get_b ()) // { dg-error "expected '\\\)' b= efore '\\\.' token" } + for (i =3D 0; i < 100; i++) + s1.get_b () +=3D 1; + +#pragma acc loop reduction(+:s1.c[1]) // { dg-error "expected '\\\)' befor= e '\\\.' token" } + for (i =3D 0; i < 100; i++) + s1.c[1] +=3D 1; + +#pragma acc loop reduction(+:s1.get_d ()[1]) // { dg-error "expected '\\\)= ' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + s1.get_d ()[1] +=3D 1; + +#pragma acc loop reduction(+:s1a[1].a) // { dg-error "expected '\\\)' befo= re '\\\[' token" } + for (i =3D 0; i < 100; i++) + s1a[1].a +=3D 1; + +#pragma acc loop reduction(+:s1a[1].get_b ()) // { dg-error "expected '\\\= )' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s1a[1].get_b () +=3D 1; + +#pragma acc loop reduction(+:s1a[1].c[1]) // { dg-error "expected '\\\)' b= efore '\\\[' token" } + for (i =3D 0; i < 100; i++) + s1a[1].c[1] +=3D 1; + +#pragma acc loop reduction(+:s1a[1].get_d ()[1]) // { dg-error "expected '= \\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s1a[1].get_d ()[1] +=3D 1; + + + // Reductions on a template struct element. + +#pragma acc loop reduction(+:s2.a) // { dg-error "expected '\\\)' before '= \\\.' token" } + for (i =3D 0; i < 100; i++) + s2.a +=3D 1; + +#pragma acc loop reduction(+:s2.get_b ()) // { dg-error "expected '\\\)' b= efore '\\\.' token" } + for (i =3D 0; i < 100; i++) + s2.get_b () +=3D 1; + +#pragma acc loop reduction(+:s2.c[1]) // { dg-error "expected '\\\)' befor= e '\\\.' token" } + for (i =3D 0; i < 100; i++) + s2.c[1] +=3D 1; + +#pragma acc loop reduction(+:s2.get_d ()[1]) // { dg-error "expected '\\\)= ' before '\\\.' token" } + for (i =3D 0; i < 100; i++) + s2.get_d ()[1] +=3D 1; + +#pragma acc loop reduction(+:s2a[1].a) // { dg-error "expected '\\\)' befo= re '\\\[' token" } + for (i =3D 0; i < 100; i++) + s2a[1].a +=3D 1; + +#pragma acc loop reduction(+:s2a[1].get_b ()) // { dg-error "expected '\\\= )' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s2a[1].get_b () +=3D 1; + +#pragma acc loop reduction(+:s2a[1].c[1]) // { dg-error "expected '\\\)' b= efore '\\\[' token" } + for (i =3D 0; i < 100; i++) + s2a[1].c[1] +=3D 1; + +#pragma acc loop reduction(+:s2a[1].get_d ()[1]) // { dg-error "expected '= \\\)' before '\\\[' token" } + for (i =3D 0; i < 100; i++) + s2a[1].get_d ()[1] +=3D 1; + + + // Reductions on arrays. + +#pragma acc loop reduction(+:a[10]) // { dg-error "expected '\\\)' before = '\\\[' token" } + for (i =3D 0; i < 100; i++) + a[10] +=3D 1; + +#pragma acc loop reduction(+:b[10]) // { dg-error "expected '\\\)' before = '\\\[' token" } + for (i =3D 0; i < 100; i++) + b[10] +=3D 1; + } +} + +int +main () +{ + test_parallel (); + test_combined (); + test_loop (); + + return 0; +} diff --git a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c b/gcc/testsuite= /gcc.dg/goacc/loop-processing-1.c index bd4c07e..1d222ab 100644 --- a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c +++ b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c @@ -15,4 +15,5 @@ void vector_1 (int *ary, int size) } } =20 -/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop 24\(1\).*= \.data_dep\.[0-9_]+ =3D \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*Head-0:.*\.= data_dep\.[0-9_]+ =3D \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*\.data_dep\.[= 0-9_]+ =3D \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data= _dep\.[0-9_]+ =3D \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.d= ata_dep\.[0-9_]+ =3D \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop = 6\(6\).*\.data_dep\.[0-9_]+ =3D \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*Head= -0:.*\.data_dep\.[0-9_]+ =3D \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*\.data_= dep\.[0-9_]+ =3D \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 1\);.*Head-1:.*= \.data_dep\.[0-9_]+ =3D \.UNIQUE \(OACC_HEAD_MARK, \.data_dep\.[0-9_]+, 1\)= ;.*\.data_dep\.[0-9_]+ =3D \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.= *Tail-1:.*\.data_dep\.[0-9_]+ =3D \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0= -9_]+, 2\);.*\.data_dep\.[0-9_]+ =3D \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9= _]+, 2\);.*Tail-0:.*\.data_dep\.[0-9_]+ =3D \.UNIQUE \(OACC_TAIL_MARK, \.da= ta_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ =3D \.UNIQUE \(OACC_JOIN, \.data= _dep\.[0-9_]+, 1\);} "oaccdevlow" } } */ +/* { dg-final { scan-tree-dump { +OpenACC loops.*Loop 0\(0\).*Loop [0-9]{2}\(1\).*\.data_dep\.[0-9_]+ =3D \.= UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*Head-0:.*\.data_dep\.[0-9_]+ =3D \.UN= IQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*\.data_dep\.[0-9_]+ =3D \.UNIQUE \(OACC= _FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data_dep\.[0-9_]+ =3D \.UNIQUE= \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ =3D \.UNI= QUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop 6\(6\).*\.data_dep\.[0-9_]= + =3D \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*Head-0:.*\.data_dep\.[0-9_]+ = =3D \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*\.data_dep\.[0-9_]+ =3D \.UNIQUE= \(OACC_FORK, \.data_dep\.[0-9_]+, 1\);.*Head-1:.*\.data_dep\.[0-9_]+ =3D \= .UNIQUE \(OACC_HEAD_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = =3D \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.*Tail-1:.*\.data_dep\.[= 0-9_]+ =3D \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 2\);.*\.data_dep= \.[0-9_]+ =3D \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 2\);.*Tail-0:.*\.d= ata_dep\.[0-9_]+ =3D \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*= \.data_dep\.[0-9_]+ =3D \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 1\);} "o= accdevlow" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90 b/gcc/= testsuite/gfortran.dg/goacc/orphan-reductions-1.f90 new file mode 100644 index 0000000..7f363d5 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90 @@ -0,0 +1,204 @@ +! Verify that gang reduction on orphan OpenACC loops reported as errors. + +subroutine s1 + implicit none + + integer, parameter :: n =3D 100 + integer :: i, sum + sum =3D 0 + + !$acc parallel reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel +end subroutine s1 + +subroutine s2 + implicit none + !$acc routine worker + + integer, parameter :: n =3D 100 + integer :: i, j, sum + sum =3D 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orph= an loop" } + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc loop reduction(+:sum) + do i =3D 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an o= rphan loop" } + do j =3D 1, n + sum =3D sum + 1 + end do + end do +end subroutine s2 + +integer function f1 () + implicit none + + integer, parameter :: n =3D 100 + integer :: i, sum + sum =3D 0 + + !$acc parallel reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel + + f1 =3D sum +end function f1 + +integer function f2 () + implicit none + !$acc routine worker + + integer, parameter :: n =3D 100 + integer :: i, j, sum + sum =3D 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orph= an loop" } + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc loop reduction(+:sum) + do i =3D 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an o= rphan loop" } + do j =3D 1, n + sum =3D sum + 1 + end do + end do + + f2 =3D sum +end function f2 + +module m +contains + subroutine s3 + implicit none + + integer, parameter :: n =3D 100 + integer :: i, sum + sum =3D 0 + + !$acc parallel reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel + end subroutine s3 + + subroutine s4 + implicit none + !$acc routine worker + + integer, parameter :: n =3D 100 + integer :: i, j, sum + sum =3D 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an or= phan loop" } + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc loop reduction(+:sum) + do i =3D 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an= orphan loop" } + do j =3D 1, n + sum =3D sum + 1 + end do + end do + end subroutine s4 + + integer function f3 () + implicit none + + integer, parameter :: n =3D 100 + integer :: i, sum + sum =3D 0 + + !$acc parallel reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i =3D 1, n + sum =3D sum + 1 + end do + !$acc end parallel + + f3 =3D sum + end function f3 + + integer function f4 () + implicit none + !$acc routine worker + + integer, parameter :: n =3D 100 + integer :: i, j, sum + sum =3D 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an or= phan loop" } + do i =3D 1, n + sum =3D sum + 1 + end do + + !$acc loop reduction(+:sum) + do i =3D 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an= orphan loop" } + do j =3D 1, n + sum =3D sum + 1 + end do + end do + + f4 =3D sum + end function f4 +end module m diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-3.c = b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-3.c new file mode 100644 index 0000000..856ef0e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-3.c @@ -0,0 +1,29 @@ +/* Check a parallel reduction which is are explicitly initialized by + the user. */ + +#include + +int +main () +{ + int n =3D 10; + float accel =3D 1.0, host =3D 1.0; + int i; + +#pragma acc parallel copyin(n) reduction(*:accel) + { + accel =3D 1.0; +#pragma acc loop gang reduction(*:accel) + for( i =3D 1; i <=3D n; i++) + { + accel *=3D 2.0; + } + } + + for (i =3D 1; i <=3D n; i++) + host *=3D 2.0; + + assert (accel =3D=3D host); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt= -2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c new file mode 100644 index 0000000..350174a --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c @@ -0,0 +1,32 @@ +#include +#include +#include + +typedef float _Complex Type; + +#define N 32 + +int +main (void) +{ + Type ary[N]; + + for (int ix =3D 0; ix < N; ix++) + ary[ix] =3D 1.0 + 1.0j; + + Type tprod =3D 1.0; + +#pragma acc parallel vector_length(32) + { +#pragma acc loop vector reduction (*:tprod) + for (int ix =3D 0; ix < N; ix++) + tprod *=3D ary[ix]; + } + + Type expected =3D 65536.0; + + if (tprod !=3D expected) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-9.f90 b/libgo= mp/testsuite/libgomp.oacc-fortran/reduction-9.f90 new file mode 100644 index 0000000..fd64d88 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/reduction-9.f90 @@ -0,0 +1,54 @@ +! Test gang reductions on dummy variables. + +! { dg-do run } + +program main + implicit none + + integer g, w, v, c + + g =3D 0 + w =3D 0 + v =3D 0 + c =3D 0 + + call reduction (g, w, v, c) + + if (g /=3D 10) call abort + if (w /=3D 10) call abort + if (v /=3D 10) call abort + if (c /=3D 100) call abort +end program main + +subroutine reduction (g, w, v, c) + implicit none + + integer g, w, v, c, i + + !$acc parallel + !$acc loop reduction(+:g) gang + do i =3D 1, 10 + g =3D g + 1 + end do + !$acc end parallel + + !$acc parallel + !$acc loop reduction(+:w) worker + do i =3D 1, 10 + w =3D w + 1 + end do + !$acc end parallel + + !$acc parallel + !$acc loop reduction(+:v) vector + do i =3D 1, 10 + v =3D v + 1 + end do + !$acc end parallel + + !$acc parallel loop reduction(+:c) gang worker vector + do i =3D 1, 100 + c =3D c + 1 + end do + !$acc end parallel loop +end subroutine reduction --MP_/.H_X2eEN/RJX2_pabxIevKR--