From 8c1ad6a29013f253137888df02b0531c08e1e09c Mon Sep 17 00:00:00 2001 From: Julian Brown Date: Tue, 26 Feb 2019 15:59:03 -0800 Subject: [PATCH] Enable firstprivate OpenACC reductions 2018-09-05 Cesar Philippidis Chung-Lin Tang 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. --- gcc/ChangeLog.omp | 6 ++ gcc/gimplify.cc | 19 ++-- gcc/testsuite/ChangeLog.omp | 5 + .../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(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/reduction-10.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/privatize-reduction-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/privatize-reduction-2.c diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 5fd096c27bc7..1f13e4a57b72 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,9 @@ +2018-09-05 Cesar Philippidis + Chung-Lin Tang + + * gimplify.cc (omp_add_variable): Enable firstprivate reduction + variables. + 2018-09-20 Cesar Philippidis * omp-low.cc (lower_oacc_head_mark): Don't mark OpenACC auto diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index a1106f28784f..d2702d153c67 100644 --- 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 f450009ab5e7..354cb06491c9 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,8 @@ +2018-09-05 Cesar Philippidis + Chung-Lin Tang + + * c-c++-common/goacc/reduction-10.c: New test. + 2018-09-20 Cesar Philippidis * 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 000000000000..579aa561479d --- /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 d06a78f321b3..afec7399353a 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,11 @@ +2018-09-05 Cesar Philippidis + Chung-Lin Tang + + * 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 * 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 000000000000..206e66fec79f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/privatize-reduction-1.c @@ -0,0 +1,41 @@ +#include +#include + +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 000000000000..0c317dcf8a6a --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/privatize-reduction-2.c @@ -0,0 +1,23 @@ +#include + +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; +} -- 2.47.2