public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-11] Enable firstprivate OpenACC reductions
@ 2021-05-13 16:09 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2021-05-13 16:09 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:a791cf334b4a6bb939b07f25e463ed528c30e4f6
commit a791cf334b4a6bb939b07f25e463ed528c30e4f6
Author: Julian Brown <julian@codesourcery.com>
Date: Tue Feb 26 15:59:03 2019 -0800
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.
Diff:
---
gcc/ChangeLog.omp | 6 ++
gcc/gimplify.c | 19 +++--
gcc/testsuite/ChangeLog.omp | 5 ++
gcc/testsuite/c-c++-common/goacc/reduction-8.c | 94 ++++++++++++++++++++++
libgomp/ChangeLog.omp | 8 ++
.../privatize-reduction-1.c | 41 ++++++++++
.../privatize-reduction-2.c | 23 ++++++
7 files changed, 189 insertions(+), 7 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index b23507ee5b9..13fade2b24b 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.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 24e4af364d5..ce09bdf35d8 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -7107,20 +7107,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
@@ -7140,9 +7147,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 2278923f564..7801b44b5f0 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 00000000000..8a0283f4ac3
--- /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 dbe4e371ee1..49ed8c68810 100644
--- 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
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;
+}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-05-13 16:09 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-13 16:09 [gcc/devel/omp/gcc-11] Enable firstprivate OpenACC reductions Kwok Yeung
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).