[gomp4, committed] Move kernels pass group before pass_fre

[gomp4, committed] Move kernels pass group before pass_fre

[Tom de Vries]

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

Hi,

this patch moves the kernels pass group to before pass_fre. Instead we 
use pass_dominator_oacc_kernels in the pass group.

This fixes an ICE while compiling the test-case included in the patch.

Committed to gomp-4_0-branch.

Thanks,
- Tom

[-- Attachment #2: 0002-Move-kernels-pass-group-before-pass_fre.patch --]
[-- Type: text/x-patch, Size: 4276 bytes --]

Move kernels pass group before pass_fre

2015-10-13  Tom de Vries  <tom@codesourcery.com>

	* tree-ssa-dom.c (pass_dominator_oacc_kernels::clone): New function.
	* passes.def: Move pass group pass_oacc_kernels to before pass_fre. Add
	pass_dominator_oacc_kernels twice in the pass_oacc_kernels pass group.

	* c-c++-common/goacc/kernels-acc-on-device-2.c: New test.
	* c-c++-common/goacc/kernels-counter-var-redundant-load.c: Update.
---
 gcc/passes.def                                     |  4 ++-
 .../c-c++-common/goacc/kernels-acc-on-device-2.c   | 37 ++++++++++++++++++++++
 .../goacc/kernels-counter-var-redundant-load.c     | 10 +++---
 gcc/tree-ssa-dom.c                                 |  1 +
 4 files changed, 47 insertions(+), 5 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device-2.c

diff --git a/gcc/passes.def b/gcc/passes.def
index bc454c0..4ed4ccd 100644
--- a/gcc/passes.def
+++ b/gcc/passes.def
@@ -86,12 +86,13 @@ along with GCC; see the file COPYING3.  If not see
 	  /* pass_build_ealias is a dummy pass that ensures that we
 	     execute TODO_rebuild_alias at this point.  */
 	  NEXT_PASS (pass_build_ealias);
-	  NEXT_PASS (pass_fre);
 	  /* Pass group that runs when there are oacc kernels in the
 	     function.  */
 	  NEXT_PASS (pass_oacc_kernels);
 	  PUSH_INSERT_PASSES_WITHIN (pass_oacc_kernels)
+	      NEXT_PASS (pass_dominator_oacc_kernels);
 	      NEXT_PASS (pass_ch_oacc_kernels);
+	      NEXT_PASS (pass_dominator_oacc_kernels);
 	      NEXT_PASS (pass_tree_loop_init);
 	      NEXT_PASS (pass_lim);
 	      NEXT_PASS (pass_copy_prop);
@@ -105,6 +106,7 @@ along with GCC; see the file COPYING3.  If not see
 	      NEXT_PASS (pass_expand_omp_ssa);
 	      NEXT_PASS (pass_tree_loop_done);
 	  POP_INSERT_PASSES ()
+	  NEXT_PASS (pass_fre);
 	  NEXT_PASS (pass_merge_phi);
           NEXT_PASS (pass_dse);
 	  NEXT_PASS (pass_cd_dce);
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device-2.c
new file mode 100644
index 0000000..2c7297b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device-2.c
@@ -0,0 +1,37 @@
+/* { dg-additional-options "-O2" } */
+
+#include "openacc.h"
+
+#define N 32
+
+void
+foo (float *a, float *b)
+{
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N])
+  {
+    int ii;
+    int on_host = acc_on_device (acc_device_X);
+
+    for (ii = 0; ii < N; ii++)
+      {
+	if (on_host)
+	  b[ii] = a[ii] + 1;
+	else
+	  b[ii] = a[ii];
+      }
+  }
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N])
+  {
+    int ii;
+    int on_host = acc_on_device (acc_device_X);
+
+    for (ii = 0; ii < N; ii++)
+      {
+	if (on_host)
+	  b[ii] = a[ii] + 2;
+	else
+	  b[ii] = a[ii];
+      }
+  }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c b/gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c
index 84dee69..c4ffc1d 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c
@@ -1,5 +1,5 @@
 /* { dg-additional-options "-O2" } */
-/* { dg-additional-options "-fdump-tree-dom_oacc_kernels" } */
+/* { dg-additional-options "-fdump-tree-dom_oacc_kernels3" } */
 
 #include <stdlib.h>
 
@@ -28,7 +28,9 @@ foo (unsigned int *c)
    _15 = .omp_data_i_10->c;
    c.1_16 = *_15;
 
-   Check that there's only one load from anonymous ssa-name (which we assume to
-   be the one to read c), and that there's no such load for ii.  */
+   Check that there are two loads from anonymous ssa-names, which we assume to
+   be:
+   - the one to read c
+   - the one to read ii after the kernels region.  */
 
-/* { dg-final { scan-tree-dump-times "(?n)\\*_\[0-9\]\[0-9\]*;$" 1 "dom_oacc_kernels" } } */
+/* { dg-final { scan-tree-dump-times "(?n)\\*_\[0-9\]\[0-9\]*;$" 2 "dom_oacc_kernels3" } } */
diff --git a/gcc/tree-ssa-dom.c b/gcc/tree-ssa-dom.c
index c7dc7b0..87f9daa 100644
--- a/gcc/tree-ssa-dom.c
+++ b/gcc/tree-ssa-dom.c
@@ -788,6 +788,7 @@ public:
   {}
 
   /* opt_pass methods: */
+  opt_pass * clone () { return new pass_dominator_oacc_kernels (m_ctxt); }
   virtual bool gate (function *) { return true; }
 
  private:
-- 
1.9.1

Re: [gomp4, committed] Move kernels pass group before pass_fre

[Tom de Vries]

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

On 13/10/15 17:59, Tom de Vries wrote:
> Hi,
>
> this patch moves the kernels pass group to before pass_fre. Instead we
> use pass_dominator_oacc_kernels in the pass group.
>

And that means we can get rid of the .omp_data_i init handling in 
tree-ssa-sccvn.c.

Committed to gomp-4_0-branch.

Thanks,
- Tom



[-- Attachment #2: 0003-Revert-.omp_data_i-init-handling-in-tree-ssa-sccvn.c.patch --]
[-- Type: text/x-patch, Size: 1143 bytes --]

Revert .omp_data_i init handling in tree-ssa-sccvn.c

2015-10-13  Tom de Vries  <tom@codesourcery.com>

	Revert:
	2015-04-21  Tom de Vries  <tom@codesourcery.com>

	* tree-ssa-sccvn.c: Include omp-low.h.
	(visit_use): Handle .omp_data_i init conservatively.
---
 gcc/tree-ssa-sccvn.c | 4 +---
 1 file changed, 1 insertion(+), 3 deletions(-)

diff --git a/gcc/tree-ssa-sccvn.c b/gcc/tree-ssa-sccvn.c
index d5964b4..5b06d29 100644
--- a/gcc/tree-ssa-sccvn.c
+++ b/gcc/tree-ssa-sccvn.c
@@ -58,7 +58,6 @@ along with GCC; see the file COPYING3.  If not see
 #include "domwalk.h"
 #include "cgraph.h"
 #include "gimple-iterator.h"
-#include "omp-low.h"
 
 /* This algorithm is based on the SCC algorithm presented by Keith
    Cooper and L. Taylor Simpson in "SCC-Based Value numbering"
@@ -3623,8 +3622,7 @@ visit_use (tree use)
     {
       if (gimple_code (stmt) == GIMPLE_PHI)
 	changed = visit_phi (stmt);
-      else if (gimple_has_volatile_ops (stmt)
-	       || gimple_stmt_omp_data_i_init_p (stmt))
+      else if (gimple_has_volatile_ops (stmt))
 	changed = defs_to_varying (stmt);
       else if (is_gimple_assign (stmt))
 	{
-- 
1.9.1