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
Move kernels pass group before pass_fre
2015-10-13 Tom de Vries <t...@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