From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1729) id 96D013938C23; Thu, 13 May 2021 16:09:16 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 96D013938C23 Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Kwok Yeung To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-11] Various OpenACC reduction enhancements - test cases X-Act-Checkin: gcc X-Git-Author: Julian Brown X-Git-Refname: refs/heads/devel/omp/gcc-11 X-Git-Oldrev: d540e4993bb2c21ad7e179a35e1181eb656974b5 X-Git-Newrev: 38cd8b95434bb7e043a51ac017cf87fd083541d1 Message-Id: <20210513160916.96D013938C23@sourceware.org> Date: Thu, 13 May 2021 16:09:16 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Thu, 13 May 2021 16:09:16 -0000 https://gcc.gnu.org/g:38cd8b95434bb7e043a51ac017cf87fd083541d1 commit 38cd8b95434bb7e043a51ac017cf87fd083541d1 Author: Julian Brown Date: Tue Feb 12 15:14:22 2019 -0800 Various OpenACC reduction enhancements - test cases 2018-12-13 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. Diff: --- gcc/testsuite/ChangeLog.omp | 11 + .../c-c++-common/goacc/orphan-reductions-1.c | 56 +++ gcc/testsuite/c-c++-common/goacc/reduction-7.c | 111 +++++ gcc/testsuite/c-c++-common/goacc/routine-4.c | 8 +- gcc/testsuite/g++.dg/goacc/reductions-1.C | 548 +++++++++++++++++++++ gcc/testsuite/gcc.dg/goacc/loop-processing-1.c | 3 +- .../gfortran.dg/goacc/orphan-reductions-1.f90 | 204 ++++++++ libgomp/ChangeLog.omp | 8 + .../libgomp.oacc-c-c++-common/par-reduction-3.c | 29 ++ .../reduction-cplx-flt-2.c | 32 ++ .../testsuite/libgomp.oacc-fortran/reduction-9.f90 | 54 ++ 11 files changed, 1059 insertions(+), 5 deletions(-) diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 5d1317d7941..798b2e261e6 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,14 @@ +2018-12-13 Cesar Philippidis + Nathan Sidwell + Julian Brown + + * 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. + 2018-06-29 Cesar Philippidis James Norris diff --git a/gcc/testsuite/c-c++-common/goacc/orphan-reductions-1.c b/gcc/testsuite/c-c++-common/goacc/orphan-reductions-1.c new file mode 100644 index 00000000000..b0bd4a7de05 --- /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 = 0; +#pragma acc loop seq reduction(+:sum) + for (i = 0; i < n; i++) + sum = sum + 1; + + return sum; +} + +#pragma acc routine gang +int +gang_reduction (int n) +{ + int i, s1 = 0, s2 = 0; +#pragma acc loop gang reduction(+:s1) /* { dg-error "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s1 = s1 + 2; + +#pragma acc loop gang reduction(+:s2) /* { dg-error "gang reduction on an orphan loop" } */ + for (i = 0; i < n; i++) + s2 = s2 + 2; + + + return s1 + s2; +} + +#pragma acc routine worker +int +worker_reduction (int n) +{ + int i, sum = 0; +#pragma acc loop worker reduction(+:sum) + for (i = 0; i < n; i++) + sum = sum + 3; + + return sum; +} + +#pragma acc routine vector +int +vector_reduction (int n) +{ + int i, sum = 0; +#pragma acc loop vector reduction(+:sum) + for (i = 0; i < n; i++) + sum = 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 00000000000..eba1d028d98 --- /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 '\\\)' before '\\\.' token" } */ + for (i = 0; i < 10; i++) + s1.a += 1; + +#pragma acc parallel reduction(+:s1.b[3]) /* { dg-error "expected '\\\)' before '\\\.' token" } */ + for (i = 0; i < 10; i++) + s1.b[3] += 1; + +#pragma acc parallel reduction(+:s2[2].a) /* { dg-error "expected '\\\)' before '\\\[' token" } */ + for (i = 0; i < 10; i++) + s2[2].a += 1; + +#pragma acc parallel reduction(+:s2[3].b[4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */ + for (i = 0; i < 10; i++) + s2[3].b[4] += 1; + +#pragma acc parallel reduction(+:z[5]) /* { dg-error "expected '\\\)' before '\\\[' token" } */ + for (i = 0; i < 10; i++) + z[5] += 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 = 0; i < 10; i++) + s1.a += 1; + +#pragma acc parallel loop reduction(+:s1.b[3]) /* { dg-error "expected '\\\)' before '\\\.' token" } */ + for (i = 0; i < 10; i++) + s1.b[3] += 1; + +#pragma acc parallel loop reduction(+:s2[2].a) /* { dg-error "expected '\\\)' before '\\\[' token" } */ + for (i = 0; i < 10; i++) + s2[2].a += 1; + +#pragma acc parallel loop reduction(+:s2[3].b[4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */ + for (i = 0; i < 10; i++) + s2[3].b[4] += 1; + +#pragma acc parallel loop reduction(+:z[5]) /* { dg-error "expected '\\\)' before '\\\[' token" } */ + for (i = 0; i < 10; i++) + z[5] += 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 = 0; i < 10; i++) + s1.a += 1; + +#pragma acc loop reduction(+:s1.b[3]) /* { dg-error "expected '\\\)' before '\\\.' token" } */ + for (i = 0; i < 10; i++) + s1.b[3] += 1; + +#pragma acc loop reduction(+:s2[2].a) /* { dg-error "expected '\\\)' before '\\\[' token" } */ + for (i = 0; i < 10; i++) + s2[2].a += 1; + +#pragma acc loop reduction(+:s2[3].b[4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */ + for (i = 0; i < 10; i++) + s2[3].b[4] += 1; + +#pragma acc loop reduction(+:z[5]) /* { dg-error "expected '\\\)' before '\\\[' token" } */ + for (i = 0; i < 10; i++) + z[5] += 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 5f2194c3f62..ad1737169f4 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-4.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-4.c @@ -23,7 +23,7 @@ void seq (void) for (int i = 0; i < 10; i++) red ++; -#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing routine" } +#pragma acc loop seq reduction (+:red) for (int i = 0; i < 10; i++) red ++; @@ -49,7 +49,7 @@ void vector (void) /* { dg-message "declared here" "1" } */ for (int i = 0; i < 10; i++) red ++; -#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing routine" } +#pragma acc loop seq reduction (+:red) for (int i = 0; i < 10; i++) red ++; @@ -75,7 +75,7 @@ void worker (void) /* { dg-message "declared here" "2" } */ for (int i = 0; i < 10; i++) red ++; -#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing routine" } +#pragma acc loop seq reduction (+:red) for (int i = 0; i < 10; i++) red ++; @@ -101,7 +101,7 @@ void gang (void) /* { dg-message "declared here" "3" } */ for (int i = 0; i < 10; i++) red ++; -#pragma acc loop gang reduction (+:red) +#pragma acc loop seq reduction (+:red) for (int i = 0; i < 10; i++) red ++; 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 00000000000..18f43f45858 --- /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 = 0; b = 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 = 0; b = 0; } + T& get_b () { return b; } + T* get_d () { return d; } +}; + +struct S1 +{ + int a, b, c[10], d[10]; + + S1 () { a = 0; b = 0; } + int& get_b () { return b; } + int* get_d () { return d; } +}; + +template +struct S2 +{ + T a, b, c[10], d[10]; + + S2 () { a = 0; b = 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 '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c1.a += 1; + +#pragma acc parallel reduction(+:c1.get_b ()) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c1.get_b () += 1; + +#pragma acc parallel reduction(+:c1.c[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c1.c[1] += 1; + +#pragma acc parallel reduction(+:c1.get_d ()[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c1.get_d ()[1] += 1; + +#pragma acc parallel reduction(+:c1a[1].a) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c1a[1].a += 1; + +#pragma acc parallel reduction(+:c1a[1].get_b ()) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c1a[1].get_b () += 1; + +#pragma acc parallel reduction(+:c1a[1].c[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c1a[1].c[1] += 1; + +#pragma acc parallel reduction(+:c1a[1].get_d ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c1a[1].get_d ()[1] += 1; + + + // Reductions on a template class member. + +#pragma acc parallel reduction(+:c2.a) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c2.a += 1; + +#pragma acc parallel reduction(+:c2.get_b ()) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c2.get_b () += 1; + +#pragma acc parallel reduction(+:c2.c[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c2.c[1] += 1; + +#pragma acc parallel reduction(+:c2.get_d ()[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c2.get_d ()[1] += 1; + + +#pragma acc parallel reduction(+:c2a[1].a) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c2a[1].a += 1; + +#pragma acc parallel reduction(+:c2a[1].get_b ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c2a[1].get_b () += 1; + +#pragma acc parallel reduction(+:c2a[1].c[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c2a[1].c[1] += 1; + +#pragma acc parallel reduction(+:c2a[1].get_d ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c2a[1].get_d ()[1] += 1; + + + // Reductions on struct element. + +#pragma acc parallel reduction(+:s1.a) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s1.a += 1; + +#pragma acc parallel reduction(+:s1.get_b ()) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s1.get_b () += 1; + +#pragma acc parallel reduction(+:s1.c[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s1.c[1] += 1; + +#pragma acc parallel reduction(+:s1.get_d ()[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s1.get_d ()[1] += 1; + +#pragma acc parallel reduction(+:s1a[1].a) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s1a[1].a += 1; + +#pragma acc parallel reduction(+:s1a[1].get_b ()) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s1a[1].get_b () += 1; + +#pragma acc parallel reduction(+:s1a[1].c[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s1a[1].c[1] += 1; + +#pragma acc parallel reduction(+:s1a[1].get_d ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s1a[1].get_d ()[1] += 1; + + + // Reductions on a template struct element. + +#pragma acc parallel reduction(+:s2.a) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s2.a += 1; + +#pragma acc parallel reduction(+:s2.get_b ()) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s2.get_b () += 1; + +#pragma acc parallel reduction(+:s2.c[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s2.c[1] += 1; + +#pragma acc parallel reduction(+:s2.get_d ()[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s2.get_d ()[1] += 1; + +#pragma acc parallel reduction(+:s2a[1].a) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s2a[1].a += 1; + +#pragma acc parallel reduction(+:s2a[1].get_b ()) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s2a[1].get_b () += 1; + +#pragma acc parallel reduction(+:s2a[1].c[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s2a[1].c[1] += 1; + +#pragma acc parallel reduction(+:s2a[1].get_d ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s2a[1].get_d ()[1] += 1; + + + // Reductions on arrays. + +#pragma acc parallel reduction(+:a[10]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + a[10] += 1; + +#pragma acc parallel reduction(+:b[10]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + b[10] += 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 = 0; i < 100; i++) + c1.a += 1; + +#pragma acc parallel loop reduction(+:c1.get_b ()) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c1.get_b () += 1; + +#pragma acc parallel loop reduction(+:c1.c[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c1.c[1] += 1; + +#pragma acc parallel loop reduction(+:c1.get_d ()[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c1.get_d ()[1] += 1; + +#pragma acc parallel loop reduction(+:c1a[1].a) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c1a[1].a += 1; + +#pragma acc parallel loop reduction(+:c1a[1].get_b ()) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c1a[1].get_b () += 1; + +#pragma acc parallel loop reduction(+:c1a[1].c[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c1a[1].c[1] += 1; + +#pragma acc parallel loop reduction(+:c1a[1].get_d ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c1a[1].get_d ()[1] += 1; + + + // Reductions on a template class member. + +#pragma acc parallel loop reduction(+:c2.a) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c2.a += 1; + +#pragma acc parallel loop reduction(+:c2.get_b ()) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c2.get_b () += 1; + +#pragma acc parallel loop reduction(+:c2.c[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c2.c[1] += 1; + +#pragma acc parallel loop reduction(+:c2.get_d ()[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c2.get_d ()[1] += 1; + + +#pragma acc parallel loop reduction(+:c2a[1].a) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c2a[1].a += 1; + +#pragma acc parallel loop reduction(+:c2a[1].get_b ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c2a[1].get_b () += 1; + +#pragma acc parallel loop reduction(+:c2a[1].c[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c2a[1].c[1] += 1; + +#pragma acc parallel loop reduction(+:c2a[1].get_d ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c2a[1].get_d ()[1] += 1; + + + // Reductions on struct element. + +#pragma acc parallel loop reduction(+:s1.a) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s1.a += 1; + +#pragma acc parallel loop reduction(+:s1.get_b ()) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s1.get_b () += 1; + +#pragma acc parallel loop reduction(+:s1.c[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s1.c[1] += 1; + +#pragma acc parallel loop reduction(+:s1.get_d ()[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s1.get_d ()[1] += 1; + +#pragma acc parallel loop reduction(+:s1a[1].a) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s1a[1].a += 1; + +#pragma acc parallel loop reduction(+:s1a[1].get_b ()) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s1a[1].get_b () += 1; + +#pragma acc parallel loop reduction(+:s1a[1].c[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s1a[1].c[1] += 1; + +#pragma acc parallel loop reduction(+:s1a[1].get_d ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s1a[1].get_d ()[1] += 1; + + + // Reductions on a template struct element. + +#pragma acc parallel loop reduction(+:s2.a) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s2.a += 1; + +#pragma acc parallel loop reduction(+:s2.get_b ()) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s2.get_b () += 1; + +#pragma acc parallel loop reduction(+:s2.c[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s2.c[1] += 1; + +#pragma acc parallel loop reduction(+:s2.get_d ()[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s2.get_d ()[1] += 1; + +#pragma acc parallel loop reduction(+:s2a[1].a) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s2a[1].a += 1; + +#pragma acc parallel loop reduction(+:s2a[1].get_b ()) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s2a[1].get_b () += 1; + +#pragma acc parallel loop reduction(+:s2a[1].c[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s2a[1].c[1] += 1; + +#pragma acc parallel loop reduction(+:s2a[1].get_d ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s2a[1].get_d ()[1] += 1; + + + // Reductions on arrays. + +#pragma acc parallel loop reduction(+:a[10]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + a[10] += 1; + +#pragma acc parallel loop reduction(+:b[10]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + b[10] += 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 = 0; i < 100; i++) + c1.a += 1; + +#pragma acc loop reduction(+:c1.get_b ()) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c1.get_b () += 1; + +#pragma acc loop reduction(+:c1.c[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c1.c[1] += 1; + +#pragma acc loop reduction(+:c1.get_d ()[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c1.get_d ()[1] += 1; + +#pragma acc loop reduction(+:c1a[1].a) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c1a[1].a += 1; + +#pragma acc loop reduction(+:c1a[1].get_b ()) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c1a[1].get_b () += 1; + +#pragma acc loop reduction(+:c1a[1].c[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c1a[1].c[1] += 1; + +#pragma acc loop reduction(+:c1a[1].get_d ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c1a[1].get_d ()[1] += 1; + + + // Reductions on a template class member. + +#pragma acc loop reduction(+:c2.a) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c2.a += 1; + +#pragma acc loop reduction(+:c2.get_b ()) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c2.get_b () += 1; + +#pragma acc loop reduction(+:c2.c[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c2.c[1] += 1; + +#pragma acc loop reduction(+:c2.get_d ()[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + c2.get_d ()[1] += 1; + + +#pragma acc loop reduction(+:c2a[1].a) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c2a[1].a += 1; + +#pragma acc loop reduction(+:c2a[1].get_b ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c2a[1].get_b () += 1; + +#pragma acc loop reduction(+:c2a[1].c[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c2a[1].c[1] += 1; + +#pragma acc loop reduction(+:c2a[1].get_d ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + c2a[1].get_d ()[1] += 1; + + + // Reductions on struct element. + +#pragma acc loop reduction(+:s1.a) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s1.a += 1; + +#pragma acc loop reduction(+:s1.get_b ()) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s1.get_b () += 1; + +#pragma acc loop reduction(+:s1.c[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s1.c[1] += 1; + +#pragma acc loop reduction(+:s1.get_d ()[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s1.get_d ()[1] += 1; + +#pragma acc loop reduction(+:s1a[1].a) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s1a[1].a += 1; + +#pragma acc loop reduction(+:s1a[1].get_b ()) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s1a[1].get_b () += 1; + +#pragma acc loop reduction(+:s1a[1].c[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s1a[1].c[1] += 1; + +#pragma acc loop reduction(+:s1a[1].get_d ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s1a[1].get_d ()[1] += 1; + + + // Reductions on a template struct element. + +#pragma acc loop reduction(+:s2.a) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s2.a += 1; + +#pragma acc loop reduction(+:s2.get_b ()) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s2.get_b () += 1; + +#pragma acc loop reduction(+:s2.c[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s2.c[1] += 1; + +#pragma acc loop reduction(+:s2.get_d ()[1]) // { dg-error "expected '\\\)' before '\\\.' token" } + for (i = 0; i < 100; i++) + s2.get_d ()[1] += 1; + +#pragma acc loop reduction(+:s2a[1].a) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s2a[1].a += 1; + +#pragma acc loop reduction(+:s2a[1].get_b ()) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s2a[1].get_b () += 1; + +#pragma acc loop reduction(+:s2a[1].c[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s2a[1].c[1] += 1; + +#pragma acc loop reduction(+:s2a[1].get_d ()[1]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + s2a[1].get_d ()[1] += 1; + + + // Reductions on arrays. + +#pragma acc loop reduction(+:a[10]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + a[10] += 1; + +#pragma acc loop reduction(+:b[10]) // { dg-error "expected '\\\)' before '\\\[' token" } + for (i = 0; i < 100; i++) + b[10] += 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 bd4c07e7d81..1d222ab3291 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) } } -/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop 24\(1\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop 6\(6\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 1\);.*Head-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.*Tail-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 2\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 2\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.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_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop 6\(6\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 1\);.*Head-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.*Tail-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 2\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 2\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 1\);} "oaccdevlow" } } */ 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 00000000000..7f363d5a5ec --- /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 = 100 + integer :: i, sum + sum = 0 + + !$acc parallel reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel +end subroutine s1 + +subroutine s2 + implicit none + !$acc routine worker + + integer, parameter :: n = 100 + integer :: i, j, sum + sum = 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do + + !$acc loop reduction(+:sum) + do i = 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do j = 1, n + sum = sum + 1 + end do + end do +end subroutine s2 + +integer function f1 () + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc parallel reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel + + f1 = sum +end function f1 + +integer function f2 () + implicit none + !$acc routine worker + + integer, parameter :: n = 100 + integer :: i, j, sum + sum = 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do + + !$acc loop reduction(+:sum) + do i = 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do j = 1, n + sum = sum + 1 + end do + end do + + f2 = sum +end function f2 + +module m +contains + subroutine s3 + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc parallel reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel + end subroutine s3 + + subroutine s4 + implicit none + !$acc routine worker + + integer, parameter :: n = 100 + integer :: i, j, sum + sum = 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do + + !$acc loop reduction(+:sum) + do i = 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do j = 1, n + sum = sum + 1 + end do + end do + end subroutine s4 + + integer function f3 () + implicit none + + integer, parameter :: n = 100 + integer :: i, sum + sum = 0 + + !$acc parallel reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel + + !$acc parallel loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + + !$acc parallel + !$acc loop gang reduction(+:sum) + do i = 1, n + sum = sum + 1 + end do + !$acc end parallel + + f3 = sum + end function f3 + + integer function f4 () + implicit none + !$acc routine worker + + integer, parameter :: n = 100 + integer :: i, j, sum + sum = 0 + + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do i = 1, n + sum = sum + 1 + end do + + !$acc loop reduction(+:sum) + do i = 1, n + !$acc loop gang reduction(+:sum) ! { dg-error "gang reduction on an orphan loop" } + do j = 1, n + sum = sum + 1 + end do + end do + + f4 = sum + end function f4 +end module m diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index df624ebfd41..f44244c6615 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,11 @@ +2018-12-13 Cesar Philippidis + Nathan Sidwell + Julian Brown + + * 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. + 2018-06-29 Cesar Philippidis James Norris 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 00000000000..856ef0e0d89 --- /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 = 10; + float accel = 1.0, host = 1.0; + int i; + +#pragma acc parallel copyin(n) reduction(*:accel) + { + accel = 1.0; +#pragma acc loop gang reduction(*:accel) + for( i = 1; i <= n; i++) + { + accel *= 2.0; + } + } + + for (i = 1; i <= n; i++) + host *= 2.0; + + assert (accel == 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 00000000000..350174a1031 --- /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 = 0; ix < N; ix++) + ary[ix] = 1.0 + 1.0j; + + Type tprod = 1.0; + +#pragma acc parallel vector_length(32) + { +#pragma acc loop vector reduction (*:tprod) + for (int ix = 0; ix < N; ix++) + tprod *= ary[ix]; + } + + Type expected = 65536.0; + + if (tprod != expected) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-9.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-9.f90 new file mode 100644 index 00000000000..fd64d88def4 --- /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 = 0 + w = 0 + v = 0 + c = 0 + + call reduction (g, w, v, c) + + if (g /= 10) call abort + if (w /= 10) call abort + if (v /= 10) call abort + if (c /= 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 = 1, 10 + g = g + 1 + end do + !$acc end parallel + + !$acc parallel + !$acc loop reduction(+:w) worker + do i = 1, 10 + w = w + 1 + end do + !$acc end parallel + + !$acc parallel + !$acc loop reduction(+:v) vector + do i = 1, 10 + v = v + 1 + end do + !$acc end parallel + + !$acc parallel loop reduction(+:c) gang worker vector + do i = 1, 100 + c = c + 1 + end do + !$acc end parallel loop +end subroutine reduction