From: Thomas Schwinge Date: Wed, 19 Oct 2016 10:19:24 +0000 (+0200) Subject: [PR tree-optimization/78024] Clear basic block flags before using BB_VISITED for... X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=a023f8c8dc9ef41f4212b964326a2d16a3f00b27;p=gcc.git [PR tree-optimization/78024] Clear basic block flags before using BB_VISITED for OpenACC loops processing gcc/ * omp-low.c (oacc_loop_discovery): Call clear_bb_flags before, and don't clear BB_VISITED after processing. gcc/testsuite/ * gcc.dg/goacc/loop-processing-1.c: New file. From-SVN: r241334 --- diff --git a/gcc/ChangeLog b/gcc/ChangeLog index d5830d552bf..4117eb34637 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,9 @@ +2016-10-19 Thomas Schwinge + + PR tree-optimization/78024 + * omp-low.c (oacc_loop_discovery): Call clear_bb_flags before, and + don't clear BB_VISITED after processing. + 2016-10-19 Richard Biener * domwalk.c (dom_walker::walk): Use RPO order. diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 77f89d5ef61..3ef796f3c24 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -19236,7 +19236,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)); @@ -19245,10 +19247,6 @@ 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; - return top; } diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index df730224983..ec875aaa5ca 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,8 @@ +2016-10-19 Thomas Schwinge + + PR tree-optimization/78024 + * gcc.dg/goacc/loop-processing-1.c: New file. + 2016-10-19 Richard Biener * gcc.dg/tree-ssa/pr61839_2.c: Fix testcase. diff --git a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c new file mode 100644 index 00000000000..619576a17ee --- /dev/null +++ b/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 (); + +void 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" } } */