Enable firstprivate OpenACC reductions

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

	gcc/
	* gimplify.cc (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.

4 files changed, 117 insertions, 7 deletions

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 5fd096c..1f13e4a 100644
--- 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.cc (omp_add_variable): Enable firstprivate reduction
+	variables.
+
 2018-09-20  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* omp-low.cc (lower_oacc_head_mark): Don't mark OpenACC auto
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 046527c..4565716 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -7381,20 +7381,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
@@ -7414,9 +7421,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 a5477c4..6e51d8f 100644
--- 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
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" } } */
+