https://gcc.gnu.org/g:38dfac029a0fddcf70bcff08adf1d68fd437662b
commit 38dfac029a0fddcf70bcff08adf1d68fd437662b
Author: Julian Brown <jul...@codesourcery.com>
Date: Tue Feb 26 15:59:03 2019 -0800
Enable firstprivate OpenACC reductions
2018-09-05 Cesar Philippidis <ce...@codesourcery.com>
Chung-Lin Tang <clt...@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:
---
gcc/ChangeLog.omp | 6 ++
gcc/gimplify.cc | 19 +++--
gcc/testsuite/ChangeLog.omp | 5 ++
gcc/testsuite/c-c++-common/goacc/reduction-10.c | 93 ++++++++++++++++++++++
libgomp/ChangeLog.omp | 8 ++
.../privatize-reduction-1.c | 41 ++++++++++
.../privatize-reduction-2.c | 23 ++++++
7 files changed, 188 insertions(+), 7 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index e2c65c4e789..1a936376617 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,9 @@
+2018-09-05 Cesar Philippidis <ce...@codesourcery.com>
+ Chung-Lin Tang <clt...@codesourcery.com>
+
+ * gimplify.cc (omp_add_variable): Enable firstprivate reduction
+ variables.
+
2018-09-20 Cesar Philippidis <ce...@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 07f835c17d9..086bcc259f5 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -7793,20 +7793,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
@@ -7826,9 +7833,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 5f40504a53e..96699d4a74f 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,8 @@
+2018-09-05 Cesar Philippidis <ce...@codesourcery.com>
+ Chung-Lin Tang <clt...@codesourcery.com>
+
+ * c-c++-common/goacc/reduction-10.c: New test.
+
2018-09-20 Cesar Philippidis <ce...@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
index 00000000000..579aa561479
--- /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 022888babdd..6c0f9951eb9 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,11 @@
+2018-09-05 Cesar Philippidis <ce...@codesourcery.com>
+ Chung-Lin Tang <clt...@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 <ce...@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
index 00000000000..206e66fec79
--- /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
index 00000000000..0c317dcf8a6
--- /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;
+}
