On Mon, Oct 17, 2016 at 11:38 AM, Thomas Schwinge <tho...@codesourcery.com> wrote: > 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?
No, just tracking in my head. > 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? Ugh. Not worse to what we do in various dwarf scanning I guess. Doesn't failure lead to a miscompile eventually? So you could formulate this as a dg-do run test with a check for the desired outcome? Richard. > 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