I've applied this patch to openacc-gcc-7-branch which privatizes reduction variables inside independent loops which are guaranteed to be assigned worker or vector partitioning during oaccdevlow. Without this patch, the inner-reduction.c test case would generate bogus results. This patch is an extension of Chung-Lin's patch posted here <https://gcc.gnu.org/ml/gcc-patches/2017-09/msg00274.html>.
The next step going forward is to teach the gimplifier to mark those reduction variables are implicitly firstprivate, instead of copy+private. Some members in the OpenACC technical committee argue that the reduction variable in inner-reduction.c should be firstprivate, not copy. Cesar
Privatize independent OpenACC reductions. 2018-01-26 Cesar Philippidis <ce...@codesourcery.com> gcc/ * gimplify.c (oacc_privatize_reduction): New function. (omp_add_variable): Use it to determine if a reduction variable needs to be privatized. libgomp/ * testsuite/libgomp.oacc-c-c++-common/inner-reduction.c: New test. diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 7a9cc241792..72ed8f1a249 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -6604,6 +6604,32 @@ omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *ctx, tree type) lang_hooks.types.omp_firstprivatize_type_sizes (ctx, type); } +/* Determine if CTX might contain any gang partitioned loops. During + oacc_dev_low, independent loops are assign gangs at the outermost + level, and vectors in the innermost. */ + +static bool +oacc_privatize_reduction (struct gimplify_omp_ctx *ctx) +{ + if (ctx == NULL) + return false; + + if (ctx->region_type != ORT_ACC) + return false; + + for (tree c = ctx->clauses; c; c = OMP_CLAUSE_CHAIN (c)) + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_SEQ: + return oacc_privatize_reduction (ctx->outer_context); + case OMP_CLAUSE_GANG: + return true; + default:; + } + + return true; +} + /* Add an entry for DECL in the OMP context CTX with FLAGS. */ static void @@ -6733,7 +6759,14 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags) } /* Set new copy map as 'private' if sure we're not gang-partitioning. */ - bool map_private = !gang && (worker || vector); + bool map_private; + + if (gang) + map_private = false; + else if (worker || vector) + map_private = true; + else + map_private = oacc_privatize_reduction (ctx->outer_context); while (outer_ctx) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/inner-reduction.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/inner-reduction.c new file mode 100644 index 00000000000..0c317dcf8a6 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/inner-reduction.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; +}