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  <cesar@codesourcery.com>
Co-Authored-By: Nathan Sidwell  <nathan@acm.org>

6 files changed, 775 insertions, 1 deletions

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 0000000..eba1d02
--- /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 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 f6e2515..d7447fd 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_]+ = \.UNIQUE \(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 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