Attached are the updated reductions tests cases. Again, these have been bootstrapped and regression tested cleanly for x86_64 with nvptx offloading. Is it OK for trunk?
Thanks, Cesar
2018-06-29 Cesar Philippidis <ce...@codesourcery.com> Nathan Sidwell <nat...@acm.org> 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. From b128e80be7cd2c81171fbd9c8b23e786bb832633 Mon Sep 17 00:00:00 2001 From: Cesar Philippidis <ce...@codesourcery.com> Date: Thu, 21 Jun 2018 11:37:56 -0700 Subject: [PATCH] Trunk reductions patches OG8 Reduction patches 4469fc4 [Fortran] Permit reductions in gfc_omp_clause_copy_ctor 704f1a2 [nxptx, OpenACC] vector reductions 8a35c89 [OpenACC] Fix a reduction bug involving GOMP_MAP_FIRSTPRIVATE_POINTER variables 16ead33 [OpenACC] Update error messages for c and c++ reductions 65dd9cf Make OpenACC orphan gang reductions errors 5d60102 [PR80547] Handle parallel reductions explicitly initialized by the user --- gcc/c/c-parser.c | 46 +- gcc/c/c-typeck.c | 8 + gcc/config/nvptx/nvptx.c | 233 +++++++- gcc/config/nvptx/nvptx.md | 7 + gcc/cp/parser.c | 27 +- gcc/cp/semantics.c | 8 + gcc/fortran/openmp.c | 12 + gcc/fortran/trans-openmp.c | 3 +- gcc/omp-general.h | 5 +- gcc/omp-low.c | 33 +- gcc/omp-offload.c | 18 + .../c-c++-common/goacc/orphan-reductions-1.c | 56 ++ .../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.dg/goacc/loop-processing-1.c | 3 +- .../gfortran.dg/goacc/orphan-reductions-1.f90 | 204 +++++++ .../par-reduction-3.c | 29 + .../reduction-cplx-flt-2.c | 32 + .../libgomp.oacc-fortran/reduction-9.f90 | 54 ++ 20 files changed, 1396 insertions(+), 49 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/orphan-reductions-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/reduction-7.c create mode 100644 gcc/testsuite/g++.dg/goacc/reductions-1.C create mode 100644 gcc/testsuite/gfortran.dg/goacc/orphan-reductions-1.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-3.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/reduction-9.f90 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 <assert.h> + +#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..245c848d509 --- /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 "invalid reduction variable" } */ + for (i = 0; i < 10; i++) + s1.a += 1; + +#pragma acc parallel reduction(+:s1.b[3]) /* { dg-error "invalid reduction variable" } */ + for (i = 0; i < 10; i++) + s1.b[3] += 1; + +#pragma acc parallel reduction(+:s2[2].a) /* { dg-error "invalid reduction variable" } */ + for (i = 0; i < 10; i++) + s2[2].a += 1; + +#pragma acc parallel reduction(+:s2[3].b[4]) /* { dg-error "invalid reduction variable" } */ + for (i = 0; i < 10; i++) + s2[3].b[4] += 1; + +#pragma acc parallel reduction(+:z[5]) /* { dg-error "invalid reduction variable" } */ + 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 "invalid reduction variable" } */ + for (i = 0; i < 10; i++) + s1.a += 1; + +#pragma acc parallel loop reduction(+:s1.b[3]) /* { dg-error "invalid reduction variable" } */ + for (i = 0; i < 10; i++) + s1.b[3] += 1; + +#pragma acc parallel loop reduction(+:s2[2].a) /* { dg-error "invalid reduction variable" } */ + for (i = 0; i < 10; i++) + s2[2].a += 1; + +#pragma acc parallel loop reduction(+:s2[3].b[4]) /* { dg-error "invalid reduction variable" } */ + for (i = 0; i < 10; i++) + s2[3].b[4] += 1; + +#pragma acc parallel loop reduction(+:z[5]) /* { dg-error "invalid reduction variable" } */ + 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 "invalid reduction variable" } */ + for (i = 0; i < 10; i++) + s1.a += 1; + +#pragma acc loop reduction(+:s1.b[3]) /* { dg-error "invalid reduction variable" } */ + for (i = 0; i < 10; i++) + s1.b[3] += 1; + +#pragma acc loop reduction(+:s2[2].a) /* { dg-error "invalid reduction variable" } */ + for (i = 0; i < 10; i++) + s2[2].a += 1; + +#pragma acc loop reduction(+:s2[3].b[4]) /* { dg-error "invalid reduction variable" } */ + for (i = 0; i < 10; i++) + s2[3].b[4] += 1; + +#pragma acc loop reduction(+:z[5]) /* { dg-error "invalid reduction variable" } */ + 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 efc4a0b95e5..91abfb5a91a 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 00000000000..6ff426a70fd --- /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 "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1.a += 1; + +#pragma acc parallel reduction(+:c1.get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1.get_b () += 1; + +#pragma acc parallel reduction(+:c1.c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1.c[1] += 1; + +#pragma acc parallel reduction(+:c1.get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1.get_d ()[1] += 1; + +#pragma acc parallel reduction(+:c1a[1].a) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1a[1].a += 1; + +#pragma acc parallel reduction(+:c1a[1].get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1a[1].get_b () += 1; + +#pragma acc parallel reduction(+:c1a[1].c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1a[1].c[1] += 1; + +#pragma acc parallel reduction(+:c1a[1].get_d ()[1]) // { dg-error "invalid reduction variable" } + 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 "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2.a += 1; + +#pragma acc parallel reduction(+:c2.get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2.get_b () += 1; + +#pragma acc parallel reduction(+:c2.c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2.c[1] += 1; + +#pragma acc parallel reduction(+:c2.get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2.get_d ()[1] += 1; + + +#pragma acc parallel reduction(+:c2a[1].a) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2a[1].a += 1; + +#pragma acc parallel reduction(+:c2a[1].get_b ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2a[1].get_b () += 1; + +#pragma acc parallel reduction(+:c2a[1].c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2a[1].c[1] += 1; + +#pragma acc parallel reduction(+:c2a[1].get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2a[1].get_d ()[1] += 1; + + + // Reductions on struct element. + +#pragma acc parallel reduction(+:s1.a) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1.a += 1; + +#pragma acc parallel reduction(+:s1.get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1.get_b () += 1; + +#pragma acc parallel reduction(+:s1.c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1.c[1] += 1; + +#pragma acc parallel reduction(+:s1.get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1.get_d ()[1] += 1; + +#pragma acc parallel reduction(+:s1a[1].a) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1a[1].a += 1; + +#pragma acc parallel reduction(+:s1a[1].get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1a[1].get_b () += 1; + +#pragma acc parallel reduction(+:s1a[1].c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1a[1].c[1] += 1; + +#pragma acc parallel reduction(+:s1a[1].get_d ()[1]) // { dg-error "invalid reduction variable" } + 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 "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2.a += 1; + +#pragma acc parallel reduction(+:s2.get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2.get_b () += 1; + +#pragma acc parallel reduction(+:s2.c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2.c[1] += 1; + +#pragma acc parallel reduction(+:s2.get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2.get_d ()[1] += 1; + +#pragma acc parallel reduction(+:s2a[1].a) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2a[1].a += 1; + +#pragma acc parallel reduction(+:s2a[1].get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2a[1].get_b () += 1; + +#pragma acc parallel reduction(+:s2a[1].c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2a[1].c[1] += 1; + +#pragma acc parallel reduction(+:s2a[1].get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2a[1].get_d ()[1] += 1; + + + // Reductions on arrays. + +#pragma acc parallel reduction(+:a[10]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + a[10] += 1; + +#pragma acc parallel reduction(+:b[10]) // { dg-error "invalid reduction variable" } + 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 "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1.a += 1; + +#pragma acc parallel loop reduction(+:c1.get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1.get_b () += 1; + +#pragma acc parallel loop reduction(+:c1.c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1.c[1] += 1; + +#pragma acc parallel loop reduction(+:c1.get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1.get_d ()[1] += 1; + +#pragma acc parallel loop reduction(+:c1a[1].a) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1a[1].a += 1; + +#pragma acc parallel loop reduction(+:c1a[1].get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1a[1].get_b () += 1; + +#pragma acc parallel loop reduction(+:c1a[1].c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1a[1].c[1] += 1; + +#pragma acc parallel loop reduction(+:c1a[1].get_d ()[1]) // { dg-error "invalid reduction variable" } + 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 "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2.a += 1; + +#pragma acc parallel loop reduction(+:c2.get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2.get_b () += 1; + +#pragma acc parallel loop reduction(+:c2.c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2.c[1] += 1; + +#pragma acc parallel loop reduction(+:c2.get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2.get_d ()[1] += 1; + + +#pragma acc parallel loop reduction(+:c2a[1].a) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2a[1].a += 1; + +#pragma acc parallel loop reduction(+:c2a[1].get_b ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2a[1].get_b () += 1; + +#pragma acc parallel loop reduction(+:c2a[1].c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2a[1].c[1] += 1; + +#pragma acc parallel loop reduction(+:c2a[1].get_d ()[1]) // { dg-error "invalid reduction variable" } + 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 "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1.a += 1; + +#pragma acc parallel loop reduction(+:s1.get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1.get_b () += 1; + +#pragma acc parallel loop reduction(+:s1.c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1.c[1] += 1; + +#pragma acc parallel loop reduction(+:s1.get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1.get_d ()[1] += 1; + +#pragma acc parallel loop reduction(+:s1a[1].a) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1a[1].a += 1; + +#pragma acc parallel loop reduction(+:s1a[1].get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1a[1].get_b () += 1; + +#pragma acc parallel loop reduction(+:s1a[1].c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1a[1].c[1] += 1; + +#pragma acc parallel loop reduction(+:s1a[1].get_d ()[1]) // { dg-error "invalid reduction variable" } + 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 "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2.a += 1; + +#pragma acc parallel loop reduction(+:s2.get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2.get_b () += 1; + +#pragma acc parallel loop reduction(+:s2.c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2.c[1] += 1; + +#pragma acc parallel loop reduction(+:s2.get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2.get_d ()[1] += 1; + +#pragma acc parallel loop reduction(+:s2a[1].a) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2a[1].a += 1; + +#pragma acc parallel loop reduction(+:s2a[1].get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2a[1].get_b () += 1; + +#pragma acc parallel loop reduction(+:s2a[1].c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2a[1].c[1] += 1; + +#pragma acc parallel loop reduction(+:s2a[1].get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2a[1].get_d ()[1] += 1; + + + // Reductions on arrays. + +#pragma acc parallel loop reduction(+:a[10]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + a[10] += 1; + +#pragma acc parallel loop reduction(+:b[10]) // { dg-error "invalid reduction variable" } + 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 "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1.a += 1; + +#pragma acc loop reduction(+:c1.get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1.get_b () += 1; + +#pragma acc loop reduction(+:c1.c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1.c[1] += 1; + +#pragma acc loop reduction(+:c1.get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1.get_d ()[1] += 1; + +#pragma acc loop reduction(+:c1a[1].a) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1a[1].a += 1; + +#pragma acc loop reduction(+:c1a[1].get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1a[1].get_b () += 1; + +#pragma acc loop reduction(+:c1a[1].c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c1a[1].c[1] += 1; + +#pragma acc loop reduction(+:c1a[1].get_d ()[1]) // { dg-error "invalid reduction variable" } + 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 "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2.a += 1; + +#pragma acc loop reduction(+:c2.get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2.get_b () += 1; + +#pragma acc loop reduction(+:c2.c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2.c[1] += 1; + +#pragma acc loop reduction(+:c2.get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2.get_d ()[1] += 1; + + +#pragma acc loop reduction(+:c2a[1].a) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2a[1].a += 1; + +#pragma acc loop reduction(+:c2a[1].get_b ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2a[1].get_b () += 1; + +#pragma acc loop reduction(+:c2a[1].c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2a[1].c[1] += 1; + +#pragma acc loop reduction(+:c2a[1].get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + c2a[1].get_d ()[1] += 1; + + + // Reductions on struct element. + +#pragma acc loop reduction(+:s1.a) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1.a += 1; + +#pragma acc loop reduction(+:s1.get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1.get_b () += 1; + +#pragma acc loop reduction(+:s1.c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1.c[1] += 1; + +#pragma acc loop reduction(+:s1.get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1.get_d ()[1] += 1; + +#pragma acc loop reduction(+:s1a[1].a) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1a[1].a += 1; + +#pragma acc loop reduction(+:s1a[1].get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1a[1].get_b () += 1; + +#pragma acc loop reduction(+:s1a[1].c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s1a[1].c[1] += 1; + +#pragma acc loop reduction(+:s1a[1].get_d ()[1]) // { dg-error "invalid reduction variable" } + 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 "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2.a += 1; + +#pragma acc loop reduction(+:s2.get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2.get_b () += 1; + +#pragma acc loop reduction(+:s2.c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2.c[1] += 1; + +#pragma acc loop reduction(+:s2.get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2.get_d ()[1] += 1; + +#pragma acc loop reduction(+:s2a[1].a) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2a[1].a += 1; + +#pragma acc loop reduction(+:s2a[1].get_b ()) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2a[1].get_b () += 1; + +#pragma acc loop reduction(+:s2a[1].c[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2a[1].c[1] += 1; + +#pragma acc loop reduction(+:s2a[1].get_d ()[1]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + s2a[1].get_d ()[1] += 1; + + + // Reductions on arrays. + +#pragma acc loop reduction(+:a[10]) // { dg-error "invalid reduction variable" } + for (i = 0; i < 100; i++) + a[10] += 1; + +#pragma acc loop reduction(+:b[10]) // { dg-error "invalid reduction variable" } + 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 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/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 <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 00000000000..350174a1031 --- /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 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 -- 2.17.1