commit 7d445a56d6db96696cec8359e58258d47fa7c9ae Author: Julian Brown Date: Wed Dec 12 11:11:03 2018 -0800 Various OpenACC reduction enhancements - test cases 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. 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 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 = 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 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 '\\\)' 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 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 = 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 ++; @@ -48,7 +48,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 ++; @@ -74,7 +74,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 ++; @@ -100,7 +100,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 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 = 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 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) } } -/* { 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 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 = 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/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 = 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 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 = 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 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 = 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