From 4566c9843f93d28f50dbe500b1b649a7731d1cd1 Mon Sep 17 00:00:00 2001 From: Cesar Philippidis Date: Tue, 26 Feb 2019 15:59:03 -0800 Subject: Enable firstprivate OpenACC reductions gcc/ChangeLog * gimplify.cc (omp_add_variable): Enable firstprivate reduction variables. gcc/testsuite/ChangeLog * c-c++-common/goacc/reduction-10.c: New test. libgomp/ChangeLog * testsuite/libgomp.oacc-c-c++-common/privatize-reduction-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/privatize-reduction-2.c: New test. Co-Authored-By: Chung-Lin Tang --- gcc/gimplify.cc | 19 +++-- gcc/testsuite/c-c++-common/goacc/reduction-10.c | 93 +++++++++++++++++++++++++ 2 files changed, 105 insertions(+), 7 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/reduction-10.c (limited to 'gcc') diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 6b45dc5..8afa993 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -8673,20 +8673,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 @@ -8706,9 +8713,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/c-c++-common/goacc/reduction-10.c b/gcc/testsuite/c-c++-common/goacc/reduction-10.c new file mode 100644 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" } } */ -- cgit v1.1