diff mbox series

[og7] Privatize independent OpenACC reductions

Message ID c2079933-fcea-61d5-9143-49df476e24a8@codesourcery.com
State New
Headers show
Series [og7] Privatize independent OpenACC reductions | expand

Commit Message

Cesar Philippidis Jan. 26, 2018, 4:38 p.m. UTC 
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
diff mbox series

Patch

Privatize independent OpenACC reductions.

2018-01-26  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.


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;
+}