This is another OpenACC reduction patch to privatize reduction variables
used inside inner acc loops. For some reason, I can't find the original
email announcement on the gcc-patches mailing list. But according to the
ChangeLog, I committed that change to og7 back on Jan 26, 2018.
I bootstrapped and regtested on x86_64/nvptx. Is it OK for trunk?
Thanks,
Cesar
>From a4753e2b40cf3d707aabd7c9d5bad7d8f9be8b6f Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <ce...@codesourcery.com>
Date: Fri, 26 Jan 2018 08:30:13 -0800
Subject: [PATCH 3/5] Privatize independent OpenACC reductions
2018-XX-YY 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.
(cherry picked from openacc-gcc-7-branch commit
330ba2316fabd0e5525c99fdacedb0bfae270244, 133f3a8fb5c)
---
gcc/gimplify.c | 35 ++++++++++++++++++-
.../inner-reduction.c | 23 ++++++++++++
2 files changed, 57 insertions(+), 1 deletion(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/inner-reduction.c
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 7dadf69b758..737a280cfe9 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6722,6 +6722,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
@@ -6851,7 +6877,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;
+}
--
2.17.1
