[gcc/devel/omp/gcc-16] Various OpenACC reduction enhancements - test cases
Sandra Loosemore via Gcc-cvs <[email protected]>
| Newsgroups | gmane.comp.gcc.cvs |
|---|---|
| Message-ID | <[email protected]> |
https://gcc.gnu.org/g:824f9ca2b0913e813a20203792d33ed2ed28f86c commit 824f9ca2b0913e813a20203792d33ed2ed28f86c Author: Julian Brown <[email protected]> Date: Sun Apr 13 23:54:16 2025 +0000 Various OpenACC reduction enhancements - test cases gcc/testsuite/ChangeLog * c-c++-common/goacc/reduction-9.c: New. * g++.dg/goacc/reductions-1.C: New. * gcc.dg/goacc/loop-processing-1.c: Update. libgomp/ChangeLog * testsuite/libgomp.oacc-c-c++-common/par-reduction-3.c: New. * testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c: New. * testsuite/libgomp.oacc-fortran/reduction-9.f90: New. Co-Authored-By: Cesar Philippidis <[email protected]> Co-Authored-By: Nathan Sidwell <[email protected]> Diff: --- gcc/testsuite/c-c++-common/goacc/reduction-9.c | 111 +++++ gcc/testsuite/g++.dg/goacc/reductions-1.C | 548 +++++++++++++++++++++ gcc/testsuite/gcc.dg/goacc/loop-processing-1.c | 2 +- .../libgomp.oacc-c-c++-common/par-reduction-3.c | 29 ++ .../reduction-cplx-flt-2.c | 32 ++ .../testsuite/libgomp.oacc-fortran/reduction-9.f90 | 54 ++ 6 files changed, 775 insertions(+), 1 deletion(-) diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-9.c b/gcc/testsuite/c-c++-common/goacc/reduction-9.c new file mode 100644 index 000000000000..eba1d028d980 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/reduction-9.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/g++.dg/goacc/reductions-1.C b/gcc/testsuite/g++.dg/goacc/reductions-1.C new file mode 100644 index 000000000000..18f43f458586 --- /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 <typename T> +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 <typename T> +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 <typename T> +void +test_parallel () +{ + int i, a[10]; + T b[10]; + C1 c1, c1a[10]; + C2<T> c2, c2a[10]; + S1 s1, s1a[10]; + S2<float> 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 <typename T> +void +test_combined () +{ + int i, a[10]; + T b[10]; + C1 c1, c1a[10]; + C2<T> c2, c2a[10]; + S1 s1, s1a[10]; + S2<float> 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 <typename T> +void +test_loop () +{ + int i, a[10]; + T b[10]; + C1 c1, c1a[10]; + C2<T> c2, c2a[10]; + S1 s1, s1a[10]; + S2<float> 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<double> (); + test_combined<long> (); + test_loop<short> (); + + 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 f6e25151e1ec..d7447fd94192 100644 --- a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c +++ b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c @@ -15,4 +15,4 @@ void vector_1 (int *ary, int size) } } -/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop 44\(1\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 68\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 68\);.*\.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_]+ = \.UNI QUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 1\);} "oaccloops" } } */ +/* { 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, 68\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 68\);.*\.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\);} "oaccloops" } } */ 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 000000000000..856ef0e0d894 --- /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 <assert.h> + +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 000000000000..350174a10318 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c @@ -0,0 +1,32 @@ +#include <complex.h> +#include <stdio.h> +#include <stdlib.h> + +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 000000000000..fd64d88def43 --- /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