Hi!
On Fri, 14 Oct 2016 13:06:59 +0200, Richard Biener <richard.guent...@gmail.com>
wrote:
> On Fri, Oct 14, 2016 at 1:00 PM, Nathan Sidwell <nat...@acm.org> wrote:
> > On 10/14/16 05:28, Richard Biener wrote:
> >
> >> The BB_VISITED flag has indetermined state at the beginning of a pass.
> >> You have to ensure it is cleared yourself.
> >
> >
> > In that case the openacc (&nvptx?) passes should be modified to clear the
> > flags at their start, rather than at their end.
The gcc/config/nvptx/nvptx.c handling seems fine -- it explicitly clears
BB_VISITED for all basic block it works on.
> Yes. But as I said, I ran into IRA ICEs (somewhere in the testsuite) when not
> cleaning up after tree-ssa-propagate.c. So somebody has to fix IRA first.
Is there a GCC PR for that, or where are you tracking such issues?
OK to commit the following? Is such a test case appropriate (which would
have caught this issue right away), in particular the dg-final
scan-tree-dump line?
commit 4e8abdfd25aa08abbad0c3fe2e9ec6182308f78c
Author: Thomas Schwinge <tho...@codesourcery.com>
Date: Mon Oct 17 11:29:43 2016 +0200
Clear basic block flags before using BB_VISITED for OpenACC loops processing
gcc/
* omp-low.c (oacc_loop_discovery): Call clear_bb_flags.
gcc/testsuite/
* gcc.dg/goacc/loop-processing-1.c: New file.
---
gcc/omp-low.c | 9 +++++----
gcc/testsuite/gcc.dg/goacc/loop-processing-1.c | 18 ++++++++++++++++++
2 files changed, 23 insertions(+), 4 deletions(-)
diff --git gcc/omp-low.c gcc/omp-low.c
index 213bf8c..5257d21 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -19340,7 +19340,9 @@ oacc_loop_sibling_nreverse (oacc_loop *loop)
static oacc_loop *
oacc_loop_discovery ()
{
- basic_block bb;
+ /* Clear basic block flags, in particular BB_VISITED which we're going to use
+ in the following. */
+ clear_bb_flags ();
oacc_loop *top = new_oacc_loop_outer (current_function_decl);
oacc_loop_discover_walk (top, ENTRY_BLOCK_PTR_FOR_FN (cfun));
@@ -19349,9 +19351,8 @@ oacc_loop_discovery ()
that diagnostics come out in an unsurprising order. */
top = oacc_loop_sibling_nreverse (top);
- /* Reset the visited flags. */
- FOR_ALL_BB_FN (bb, cfun)
- bb->flags &= ~BB_VISITED;
+ /* Clear basic block flags again, as otherwise IRA will explode later on. */
+ clear_bb_flags ();
return top;
}
diff --git gcc/testsuite/gcc.dg/goacc/loop-processing-1.c
gcc/testsuite/gcc.dg/goacc/loop-processing-1.c
new file mode 100644
index 0000000..2f0b3a2
--- /dev/null
+++ gcc/testsuite/gcc.dg/goacc/loop-processing-1.c
@@ -0,0 +1,18 @@
+/* Make sure that OpenACC loop processing happens. */
+/* { dg-additional-options "-O2 -fdump-tree-oaccdevlow" } */
+
+extern int place ();
+
+int vector_1 (int *ary, int size)
+{
+#pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size])
firstprivate (size)
+ {
+#pragma acc loop gang
+ for (int jx = 0; jx < 1; jx++)
+#pragma acc loop auto
+ for (int ix = 0; ix < size; ix++)
+ ary[ix] = place ();
+ }
+}
+
+/* { dg-final { scan-tree-dump "OpenACC loops.*Loop 0\\\(0\\\).*Loop
14\\\(1\\\).*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, 0, 1,
20\\\);.*Head-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, 0, 1,
20\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_FORK,
\\\.data_dep\\\.\[0-9_\]+, 0\\\);.*Tail-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE
\\\(OACC_TAIL_MARK, \\\.data_dep\\\.\[0-9_\]+,
1\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_JOIN,
\\\.data_dep\\\.\[0-9_\]+, 0\\\);.*Loop 6\\\(4\\\).*\\\.data_dep\\\.\[0-9_\]+ =
UNIQUE \\\(OACC_HEAD_MARK, 0, 1, 6\\\);.*Head-0:.*\\\.data_dep\\\.\[0-9_\]+ =
UNIQUE \\\(OACC_HEAD_MARK, 0, 1, 6\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE
\\\(OACC_FORK, \\\.data_dep\\\.\[0-9_\]+,
2\\\);.*Tail-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_TAIL_MARK,
\\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE
\\\(OACC_JOIN, \\\.data_dep\\\.\[0-9_\]+, 2\\\);" "oaccdevlow" } } */
Grüße
Thomas
