Add pass_dominator_oacc_kernels

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

	* passes.def: Add pass_dominator_oacc_kernels to pass group pass_oacc_kernels.
	Add pass_tree_loop_done before, and pass_tree_loop_init after.
	* tree-pass.h (make_pass_dominator_oacc_kernels): Declare.
	* tree-ssa-dom.c (class dominator_base): New class.  Factor out of ...
	(class pass_dominator): ... here.
	(pass_dominator_oacc_kernels): New pass.
	(make_pass_dominator_oacc_kernels): New function.

	* c-c++-common/goacc/kernels-counter-var-redundant-load.c: New test.
---
 gcc/passes.def                                     |   3 +
 .../goacc/kernels-counter-var-redundant-load.c     |  34 ++++++
 gcc/tree-pass.h                                    |   1 +
 gcc/tree-ssa-dom.c                                 | 117 +++++++++++++++++----
 4 files changed, 134 insertions(+), 21 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c

diff --git a/gcc/passes.def b/gcc/passes.def
index 0498a8b..bc454c0 100644
--- a/gcc/passes.def
+++ b/gcc/passes.def
@@ -98,6 +98,9 @@ along with GCC; see the file COPYING3.  If not see
 	      NEXT_PASS (pass_lim);
 	      NEXT_PASS (pass_copy_prop);
 	      NEXT_PASS (pass_scev_cprop);
+	      NEXT_PASS (pass_tree_loop_done);
+	      NEXT_PASS (pass_dominator_oacc_kernels);
+	      NEXT_PASS (pass_tree_loop_init);
       	      NEXT_PASS (pass_parallelize_loops_oacc_kernels);
 	      NEXT_PASS (pass_expand_omp_ssa);
 	      NEXT_PASS (pass_tree_loop_done);
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
new file mode 100644
index 0000000..84dee69
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-counter-var-redundant-load.c
@@ -0,0 +1,34 @@
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fdump-tree-dom_oacc_kernels" } */
+
+#include <stdlib.h>
+
+#define N (1024 * 512)
+#define COUNTERTYPE unsigned int
+
+COUNTERTYPE
+foo (unsigned int *c)
+{
+  COUNTERTYPE ii;
+
+#pragma acc kernels copyout (c[0:N])
+  {
+    for (ii = 0; ii < N; ii++)
+      c[ii] = 1;
+  }
+
+  return ii;
+}
+
+/* We're expecting:
+
+   .omp_data_i_10 = &.omp_data_arr.3;
+   _11 = .omp_data_i_10->ii;
+   *_11 = 0;
+   _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.  */
+
+/* { dg-final { scan-tree-dump-times "(?n)\\*_\[0-9\]\[0-9\]*;$" 1 "dom_oacc_kernels" } } */
diff --git a/gcc/tree-pass.h b/gcc/tree-pass.h
index 52ba3e5..15c8bf6 100644
--- a/gcc/tree-pass.h
+++ b/gcc/tree-pass.h
@@ -392,6 +392,7 @@ extern gimple_opt_pass *make_pass_build_ssa (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_build_alias (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_build_ealias (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_dominator (gcc::context *ctxt);
+extern gimple_opt_pass *make_pass_dominator_oacc_kernels (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_dce (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_cd_dce (gcc::context *ctxt);
 extern gimple_opt_pass *make_pass_call_cdce (gcc::context *ctxt);
diff --git a/gcc/tree-ssa-dom.c b/gcc/tree-ssa-dom.c
index 573e6fc..c7dc7b0 100644
--- a/gcc/tree-ssa-dom.c
+++ b/gcc/tree-ssa-dom.c
@@ -45,6 +45,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "gimplify.h"
 #include "tree-cfgcleanup.h"
 #include "cfgcleanup.h"
+#include "omp-low.h"
 
 /* This file implements optimizations on the dominator tree.  */
 
@@ -526,6 +527,31 @@ private:
 
 namespace {
 
+class dominator_base : public gimple_opt_pass
+{
+ protected:
+  dominator_base (pass_data data, gcc::context *ctxt)
+    : gimple_opt_pass (data, ctxt)
+  {}
+
+  unsigned int execute (function *);
+
+  /* Return true if pass should perform jump threading.  */
+  virtual bool jump_threading_p (void) { return !sese_mode_p (); }
+
+  /* Return true if pass should visit a series of seses rather than the whole
+     dominator tree.  */
+  virtual bool sese_mode_p (void) { return false; }
+
+  /* In sese mode, return true if there's another sese to visit.  Return the
+     sese to visit in SESE_ENTRY and SESE_EXIT.  If a stmt in the sese should
+     not be optimized, return it in SKIP_STMT.  */
+  virtual bool get_sese (basic_block *sese_entry ATTRIBUTE_UNUSED,
+			 basic_block *sese_exit ATTRIBUTE_UNUSED,
+			 gimple **skip_stmt ATTRIBUTE_UNUSED)
+    { gcc_unreachable (); }
+}; // class dominator_base
+
 const pass_data pass_data_dominator =
 {
   GIMPLE_PASS, /* type */
@@ -539,38 +565,20 @@ const pass_data pass_data_dominator =
   ( TODO_cleanup_cfg | TODO_update_ssa ), /* todo_flags_finish */
 };
 
-class pass_dominator : public gimple_opt_pass
+class pass_dominator : public dominator_base
 {
 public:
   pass_dominator (gcc::context *ctxt)
-    : gimple_opt_pass (pass_data_dominator, ctxt)
+    : dominator_base (pass_data_dominator, ctxt)
   {}
 
   /* opt_pass methods: */
   opt_pass * clone () { return new pass_dominator (m_ctxt); }
   virtual bool gate (function *) { return flag_tree_dom != 0; }
-  virtual unsigned int execute (function *);
-
- protected:
-  /* Return true if pass should perform jump threading.  */
-  virtual bool jump_threading_p (void) { return !sese_mode_p (); }
-
-  /* Return true if pass should visit a series of seses rather than the whole
-     dominator tree.  */
-  virtual bool sese_mode_p (void) { return false; }
-
-  /* In sese mode, return true if there's another sese to visit.  Return the
-     sese to visit in SESE_ENTRY and SESE_EXIT.  If a stmt in the sese should
-     not be optimized, return it in SKIP_STMT.  */
-  virtual bool get_sese (basic_block *sese_entry ATTRIBUTE_UNUSED,
-			 basic_block *sese_exit ATTRIBUTE_UNUSED,
-			 gimple **skip_stmt ATTRIBUTE_UNUSED)
-    { gcc_unreachable (); }
-
 }; // class pass_dominator
 
 unsigned int
-pass_dominator::execute (function *fun)
+dominator_base::execute (function *fun)
 {
   memset (&opt_stats, 0, sizeof (opt_stats));
 
@@ -759,6 +767,68 @@ pass_dominator::execute (function *fun)
   return 0;
 }
 
+const pass_data pass_data_dominator_oacc_kernels =
+{
+  GIMPLE_PASS, /* type */
+  "dom_oacc_kernels", /* name */
+  OPTGROUP_NONE, /* optinfo_flags */
+  TV_TREE_SSA_DOMINATOR_OPTS, /* tv_id */
+  ( PROP_cfg | PROP_ssa ), /* properties_required */
+  0, /* properties_provided */
+  0, /* properties_destroyed */
+  0, /* todo_flags_start */
+  ( TODO_cleanup_cfg | TODO_update_ssa ), /* todo_flags_finish */
+};
+
+class pass_dominator_oacc_kernels : public dominator_base
+{
+public:
+  pass_dominator_oacc_kernels (gcc::context *ctxt)
+    : dominator_base (pass_data_dominator_oacc_kernels, ctxt), m_regions (NULL)
+  {}
+
+  /* opt_pass methods: */
+  virtual bool gate (function *) { return true; }
+
+ private:
+  bitmap m_regions;
+
+protected:
+  /* dominator_base methods: */
+  virtual bool sese_mode_p (void) { return true; }
+  virtual bool get_sese (basic_block *sese_entry, basic_block *sese_exit,
+			 gimple **skip_stmt)
+  {
+    if (m_regions == NULL)
+      {
+	m_regions = BITMAP_ALLOC (NULL);
+	basic_block bb;
+	FOR_EACH_BB_FN (bb, cfun)
+	  if (oacc_kernels_region_entry_p (bb, NULL))
+	    bitmap_set_bit (m_regions, bb->index);
+      }
+
+    if (bitmap_empty_p (m_regions))
+      {
+	BITMAP_FREE (m_regions);
+	return false;
+      }
+
+    unsigned int index = bitmap_first_set_bit (m_regions);
+    bitmap_clear_bit (m_regions, index);
+
+    *sese_entry = BASIC_BLOCK_FOR_FN (cfun, index);
+    *sese_exit = get_oacc_kernels_region_exit (*sese_entry);
+
+    tree omp_data_i = get_omp_data_i (single_pred (*sese_entry));
+    if (omp_data_i != NULL_TREE)
+      *skip_stmt = SSA_NAME_DEF_STMT (omp_data_i);
+
+    return true;
+  }
+
+}; // class pass_dominator_oacc_kernels
+
 } // anon namespace
 
 gimple_opt_pass *
@@ -767,6 +837,11 @@ make_pass_dominator (gcc::context *ctxt)
   return new pass_dominator (ctxt);
 }
 
+gimple_opt_pass *
+make_pass_dominator_oacc_kernels (gcc::context *ctxt)
+{
+  return new pass_dominator_oacc_kernels (ctxt);
+}
 
 /* Given a conditional statement CONDSTMT, convert the
    condition to a canonical form.  */
-- 
1.9.1

