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
@@ -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);
new file mode 100644
@@ -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" } } */
@@ -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);
@@ -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