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-10.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.

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 5fd096c27bc76722ec752c7bc83ea3b144fa1163..1f13e4a57b72ba61a3b97df9dee1ef434cdad358 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.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 a1106f28784fb40432c53a8a6c588d8aae8ca5d0..d2702d153c67331f616b77c2011e545ab1310bb8 100644 (file)
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -7449,20 +7449,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
@@ -7482,9 +7489,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 f450009ab5e764f6909ec175f538b5e9587ab081..354cb06491c92ac5c27459895bc5ddf751edc23c 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-10.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-10.c b/gcc/testsuite/c-c++-common/goacc/reduction-10.c
new file mode 100644 (file)
index 0000000..579aa56
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-10.c
@@ -0,0 +1,93 @@
+/* { 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 d06a78f321b3846e33cfe279043bad2472418e09..afec7399353a147014678bd7f418d27f6172cf93 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;
+}