Enable firstprivate OpenACC reductions

2018-09-05  Cesar Philippidis  <cesar@codesourcery.com>
	    Chung-Lin Tang  <cltang@codesourcery.com>

	gcc/
	* gimplify.c (omp_add_variable): Enable firstprivate reduction
	variables.

	gcc/testsuite/
	* c-c++-common/goacc/reduction-8.c: New test.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/privatize-reduction-1.c: New
	test.
	* testsuite/libgomp.oacc-c-c++-common/privatize-reduction-2.c: New
	test.

(cherry picked from openacc-gcc-9-branch commit
1add5113c32e7b8ea8536aee877f901a4a6eb5fc)

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 45ed71e87736bd9d767a713118be49fd0bb6946c..0e76faedc3f5ed58ddc494d0070465956b1b6dae 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,9 @@
+2018-09-05  Cesar Philippidis  <cesar@codesourcery.com>
+           Chung-Lin Tang  <cltang@codesourcery.com>
+
+       * gimplify.c (omp_add_variable): Enable firstprivate reduction
+       variables.
+
 2018-09-20  Cesar Philippidis  <cesar@codesourcery.com>
 
        * omp-low.c (lower_oacc_head_mark): Don't mark OpenACC auto

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 3034a252e7ef1d5a10c653f4baf2e5be3481c19e..e253c2b0c27010a945a049f288e88973bcaadf58 100644 (file)
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6941,20 +6941,27 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
   else
     splay_tree_insert (ctx->variables, (splay_tree_key)decl, flags);
 
-  /* For reductions clauses in OpenACC loop directives, by default create a
-     copy clause on the enclosing parallel construct for carrying back the
-     results.  */
+  /* For OpenACC loop directives, when a reduction clause is placed on
+     the outermost acc loop within an acc parallel or kernels
+     construct, it must have an implied copy data mapping. E.g.
+
+       #pragma acc parallel
+        {
+          #pragma acc loop reduction (+:sum)
+
+     a copy clause for sum should be added on the enclosing parallel
+     construct for carrying back the results.  */
   if (ctx->region_type == ORT_ACC && (flags & GOVD_REDUCTION))
     {
       struct gimplify_omp_ctx *outer_ctx = ctx->outer_context;
-      while (outer_ctx)
+      if (outer_ctx)
        {
          n = splay_tree_lookup (outer_ctx->variables, (splay_tree_key)decl);
          if (n != NULL)
            {
              /* Ignore local variables and explicitly declared clauses.  */
              if (n->value & (GOVD_LOCAL | GOVD_EXPLICIT))
-               break;
+               ;
              else if (outer_ctx->region_type == ORT_ACC_KERNELS)
                {
                  /* According to the OpenACC spec, such a reduction variable
@@ -6974,9 +6981,7 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
            {
              splay_tree_insert (outer_ctx->variables, (splay_tree_key)decl,
                                 GOVD_MAP | GOVD_SEEN);
-             break;
            }
-         outer_ctx = outer_ctx->outer_context;
        }
     }
 }

diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 803fa3baa9e4f9329af66f6887eaa0503d2862e7..7f283b6063f03b5bef9cf8b138269efff476a55c 100644 (file)
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,8 @@
+2018-09-05  Cesar Philippidis  <cesar@codesourcery.com>
+           Chung-Lin Tang  <cltang@codesourcery.com>
+
+       * c-c++-common/goacc/reduction-8.c: New test.
+
 2018-09-20  Cesar Philippidis  <cesar@codesourcery.com>
 
        * c-c++-common/goacc/loop-auto-1.c: Adjust test case to conform to

diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-8.c b/gcc/testsuite/c-c++-common/goacc/reduction-8.c
new file mode 100644 (file)
index 0000000..8a0283f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-8.c
@@ -0,0 +1,94 @@
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define n 1000
+
+int
+main(void)
+{
+  int i, j;
+  int result, array[n];
+
+#pragma acc parallel loop reduction (+:result)
+  for (i = 0; i < n; i++)
+    result ++;
+
+#pragma acc parallel
+#pragma acc loop reduction (+:result)
+  for (i = 0; i < n; i++)
+    result ++;
+
+#pragma acc parallel
+#pragma acc loop
+  for (i = 0; i < n; i++)
+    {
+      result = i;
+
+#pragma acc loop reduction(+:result)
+      for (j = 0; j < n; j++)
+       result ++;
+
+      array[i] = result;
+    }
+
+#pragma acc parallel
+#pragma acc loop
+  for (i = 0; i < n; i++)
+    {
+      result = i;
+
+#pragma acc loop worker vector reduction(+:result)
+      for (j = 0; j < n; j++)
+       result ++;
+
+      array[i] = result;
+    }
+
+#pragma acc parallel
+#pragma acc loop // { dg-warning "insufficient partitioning" }
+  for (i = 0; i < n; i++)
+    {
+      result = i;
+
+#pragma acc loop gang reduction(+:result)
+      for (j = 0; j < n; j++)
+       result ++;
+
+      array[i] = result;
+    }
+
+#pragma acc parallel copy(result)
+#pragma acc loop // { dg-warning "insufficient partitioning" }
+  for (i = 0; i < n; i++)
+    {
+      result = i;
+
+#pragma acc loop gang reduction(+:result)
+      for (j = 0; j < n; j++)
+       result ++;
+
+      array[i] = result;
+    }
+  
+#pragma acc kernels
+#pragma acc loop
+  for (i = 0; i < n; i++)
+    {
+      result = i;
+
+#pragma acc loop reduction(+:result)
+      for (j = 0; j < n; j++)
+       result ++;
+
+      array[i] = result;
+    }
+
+  return 0;
+}
+
+/* Check that default copy maps are generated for loop reductions.  */
+/* { dg-final { scan-tree-dump-times "reduction..:result. map.tofrom:result .len: 4.." 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "oacc_parallel map.tofrom:result .len: 4.." 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000.. firstprivate.result." 3 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map.tofrom:result .len: 4.. map.tofrom:array .len: 4000.." 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000.. map.force_tofrom:result .len: 4.." 1 "gimple" } } */
+

diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 82674a8ee057216207b72cd7b1045abebd7e1980..382126b86bba66d655787094f4dfa873f0a9fd52 100644 (file)
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,11 @@
+2018-09-05  Cesar Philippidis  <cesar@codesourcery.com>
+           Chung-Lin Tang  <cltang@codesourcery.com>
+
+       * testsuite/libgomp.oacc-c-c++-common/privatize-reduction-1.c: New
+       test.
+       * testsuite/libgomp.oacc-c-c++-common/privatize-reduction-2.c: New
+       test.
+
 2018-09-20  Cesar Philippidis  <cesar@codesourcery.com>
 
        * testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Adjust test case

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/privatize-reduction-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/privatize-reduction-1.c
new file mode 100644 (file)
index 0000000..206e66f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/privatize-reduction-1.c
@@ -0,0 +1,41 @@
+#include <stdio.h>
+#include <stdlib.h>
+
+int
+main (int argc, char *argv[])
+{
+#define N 100
+  int n = N;
+  int i, j, tmp;
+  int input[N*N], output[N], houtput[N];
+
+  for (i = 0; i < n * n; i++)
+    input[i] = i;
+
+  for (i = 0; i < n; i++)
+    {
+      tmp = 0;
+      for (j = 0; j < n; j++)
+       tmp += input[i * n + j];
+      houtput[i] = tmp;
+    }
+
+  #pragma acc parallel loop gang
+  for (i = 0; i < n; i++)
+    {
+      tmp = 0;
+
+      #pragma acc loop worker reduction(+:tmp)
+      for (j = 0; j < n; j++)
+       tmp += input[i * n + j];
+
+      output[i] = tmp;
+    }
+
+  /* Test if every worker-level reduction had correct private result.  */
+  for (i = 0; i < n; i++)
+    if (houtput[i] != output[i])
+      abort ();
+
+  return 0;
+}

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/privatize-reduction-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/privatize-reduction-2.c
new file mode 100644 (file)
index 0000000..0c317dc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/privatize-reduction-2.c
@@ -0,0 +1,23 @@
+#include <assert.h>
+
+int
+main ()
+{
+  const int n = 1000;
+  int i, j, temp, a[n];
+
+#pragma acc parallel loop
+  for (i = 0; i < n; i++)
+    {
+      temp = i;
+#pragma acc loop reduction (+:temp)
+      for (j = 0; j < n; j++)
+       temp ++;
+      a[i] = temp;
+    }
+
+  for (i = 0; i < n; i++)
+    assert (a[i] == i+n);
+
+  return 0;
+}