This is the mail archive of the
gcc@gcc.gnu.org
mailing list for the GCC project.
[PR tree-optimization/78024] Clear basic block flags before using BB_VISITED for OpenACC loops processing
Hi!
On Wed, 19 Oct 2016 12:07:13 +0200, Richard Biener <richard.guenther@gmail.com> wrote:
> On Tue, Oct 18, 2016 at 9:52 PM, Thomas Schwinge
> <thomas@codesourcery.com> wrote:
> > can I at
> > least commit the OpenACC loops processing fix? Here is the latest
> > version, simplified after your r241296 IRA vs. BB_VISITED fixes:
>
> Sure, I considered that approved already
I don't see such an approval anywhere?
> (it's even obvious).
I've been told to be very careful in GCC with declaring something
"obvious". (Though, I will be happy to increase my autonomy in declaring
patches "obvious"/"ready for commit". This will potentially save me from
wasting hours when keeping patches up-to-date, while repeatedly begging
for review, and so on.)
Anyway, without changes now committed in r241334:
commit b3e3b38b4b1ac65df70d98f7e7557a27947948c1
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Oct 19 10:19:24 2016 +0000
[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.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@241334 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 6 ++++++
gcc/omp-low.c | 8 +++-----
gcc/testsuite/ChangeLog | 5 +++++
gcc/testsuite/gcc.dg/goacc/loop-processing-1.c | 18 ++++++++++++++++++
4 files changed, 32 insertions(+), 5 deletions(-)
diff --git gcc/ChangeLog gcc/ChangeLog
index d5830d5..4117eb3 100644
--- gcc/ChangeLog
+++ gcc/ChangeLog
@@ -1,3 +1,9 @@
+2016-10-19 Thomas Schwinge <thomas@codesourcery.com>
+
+ 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 <rguenther@suse.de>
* domwalk.c (dom_walker::walk): Use RPO order.
diff --git gcc/omp-low.c gcc/omp-low.c
index 77f89d5..3ef796f 100644
--- gcc/omp-low.c
+++ 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 gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog
index df73022..ec875aa 100644
--- gcc/testsuite/ChangeLog
+++ gcc/testsuite/ChangeLog
@@ -1,3 +1,8 @@
+2016-10-19 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR tree-optimization/78024
+ * gcc.dg/goacc/loop-processing-1.c: New file.
+
2016-10-19 Richard Biener <rguenther@suse.de>
* gcc.dg/tree-ssa/pr61839_2.c: Fix testcase.
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..619576a
--- /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 ();
+
+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" } } */
Grüße
Thomas