Various OpenACC reduction enhancements - test cases

2018-12-13  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.

(cherry picked from openacc-gcc-9-branch commit
e3cdb913b543ebd719dc8fe33f92f504c141d9f6)

diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index e379c43cffe579b73eca7c72a12356c9e498a5e2..9f9086f82e81037245d78a04fcb845b4161d956c 100644 (file)
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,14 @@
+2018-12-13  Cesar Philippidis  <cesar@codesourcery.com>
+            Nathan Sidwell  <nathan@acm.org>
+            Julian Brown  <julian@codesourcery.com>
+
+        * 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.
+
 2018-06-29  Cesar Philippidis  <cesar@codesourcery.com>
            James Norris  <jnorris@codesourcery.com>
 

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 (file)
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 (file)
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 5f2194c3f6238dae9cd708d932914b15cb4213bd..ad1737169f4f27177907d57206426816fd781060 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/routine-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-4.c
@@ -23,7 +23,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 ++;
 
@@ -49,7 +49,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 ++;
 
@@ -75,7 +75,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 ++;
 
@@ -101,7 +101,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 (file)
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 bd4c07e7d81cf4334912bea1711a983b32a811e2..1d222ab329161522231c60b05ed0274776d6d493 100644 (file)
--- 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 (file)
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/ChangeLog.omp b/libgomp/ChangeLog.omp
index 9f5e057350c6e4533be0ff53f96d0776c46d495d..f810768a4c1c31fc1c52d4624687cc99a2d17303 100644 (file)
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,11 @@
+2018-12-13  Cesar Philippidis  <cesar@codesourcery.com>
+            Nathan Sidwell  <nathan@acm.org>
+            Julian Brown  <julian@codesourcery.com>
+
+        * 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.
+
 2018-06-29  Cesar Philippidis  <cesar@codesourcery.com>
            James Norris  <jnorris@codesourcery.com>
 

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 (file)
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 (file)
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 (file)
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