[PATCH] Privatize independent OpenACC reductions

[Cesar Philippidis <cesar at codesourcery dot com>]

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 <cesar@codesourcery.com>
Date: Fri, 26 Jan 2018 08:30:13 -0800
Subject: [PATCH 3/5] Privatize independent OpenACC reductions

2018-XX-YY  Cesar Philippidis  <cesar@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