This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [patch] various OpenACC reduction enhancements - test cases
- From: Julian Brown <julian at codesourcery dot com>
- To: Jakub Jelinek <jakub at redhat dot com>
- Cc: Cesar Philippidis <cesar at codesourcery dot com>, Thomas Schwinge <thomas at codesourcery dot com>, "gcc-patches at gcc dot gnu dot org" <gcc-patches at gcc dot gnu dot org>, Tom de Vries <tdevries at suse dot de>
- Date: Thu, 13 Dec 2018 14:13:58 +0000
- Subject: Re: [patch] various OpenACC reduction enhancements - test cases
- References: <ae587152-072d-52b0-0129-99de0ed40d03@codesourcery.com> <afa69e6b-1b09-63f2-556f-e167f8b645fe@codesourcery.com> <20181204125933.GM12380@tucnak>
On Tue, 4 Dec 2018 13:59:33 +0100
Jakub Jelinek <jakub@redhat.com> wrote:
> On Fri, Jun 29, 2018 at 11:23:21AM -0700, Cesar Philippidis wrote:
> > Attached are the updated reductions tests cases. Again, these have
> > been bootstrapped and regression tested cleanly for x86_64 with
> > nvptx offloading. Is it OK for trunk?
>
> If Thomas is ok with this, it is ok for trunk.
Here's a new version to go with the FE patch posted here:
https://gcc.gnu.org/ml/gcc-patches/2018-12/msg00930.html
Thanks,
Julian
ChangeLog
2018-xx-xx Cesar Philippidis <cesar@codesourcery.com>
Nathan Sidwell <nathan@acm.org>
Julian Brown <julian@codesourcery.com>
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.
commit 7d445a56d6db96696cec8359e58258d47fa7c9ae
Author: Julian Brown <julian@codesourcery.com>
Date: Wed Dec 12 11:11:03 2018 -0800
Various OpenACC reduction enhancements - test cases
2018-xx-xx Cesar Philippidis <cesar@codesourcery.com>
Nathan Sidwell <nathan@acm.org>
Julian Brown <julian@codesourcery.com>
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 <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 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 <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 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 <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 0000000..350174a
--- /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 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