[gomp4] Handle parallel reductions explicitly initialized by the user

[gomp4] Handle parallel reductions explicitly initialized by the user

[Cesar Philippidis]

[-- Attachment #1: Type: text/plain, Size: 1267 bytes --]

This patch resolves an ICE inside the nvptx BE involving reduction
variables which are initialized by the user. E.g.

  #pragma acc parallel reduction(+:var)
  {
    var = 1;

    #pragma acc loop reduction(+:var)
    ...

Currently, the nvptx BE expects the internal function GOACC_REDUCTION to
have a LHS argument containing the reduction variable to be initialized,
however in this case that would be overly aggressive. So when the user
initializes the reduction variable explicitly to an immediate constant,
the optimizer will propagate that value to the reduction initializer
associated with inner the acc loop directly and remove the LHS for the
outer parallel reduction initializer. This shouldn't be a problem here
because during omp lowering, reduction variables are initialized just
outside of the loop. So even if the loop is a empty, the outer parallel
reduction will still be initialized properly.

The other approach I supposed would be to inhibit constant propagation
on reduction on those GOACC internal functions. But that is probably too
restrictive.

I applied this patch to gomp-4_0-branch. Bernd, I believe that trunk may
have a similar problem. If does, is this OK for trunk? I'll hopefully
start working on trunk again in a week or so.

Cesar


[-- Attachment #2: gomp4-parallel_reduction_ice.diff --]
[-- Type: text/x-patch, Size: 1542 bytes --]

2017-05-12  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* config/nvptx/nvptx.c (nvptx_goacc_reduction_init): Don't update
	LHS if it doesn't exist.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/par-reduction-3.c: New test.


diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index f790728..bd1236b 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -4948,7 +4948,11 @@ nvptx_goacc_reduction_init (gcall *call)
 	    init = var;
 	}
 
-      gimplify_assign (lhs, init, &seq);
+      /* The LHS may be NULL if a reduction variable on a parallel
+	 construct is initialized to some constant inside the parallel
+	 region.  */
+      if (lhs)
+	gimplify_assign (lhs, init, &seq);
     }
 
   pop_gimplify_context (NULL);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-3.c
new file mode 100644
index 0000000..6510143
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/par-reduction-3.c
@@ -0,0 +1,31 @@
+/* Check a parallel reduction which is are explicitly initialized by
+   the user.  */
+
+/* { dg-do run } */
+
+#include <assert.h>
+
+int
+main ()
+{
+  int n = 10;
+  float accel = 1.0, host = 1.0;
+  int i;
+
+#pragma acc parallel copyin(n) reduction(*:accel)
+  {
+    accel = 1.0;
+#pragma acc loop gang reduction(*:accel)
+    for( i = 1; i <= n; i++)
+      {
+	accel *= 2.0;
+      }
+  }
+
+  for (i = 1; i <= n; i++)
+    host *= 2.0;
+
+  assert (accel == host);
+
+  return 0;
+}