[PR tree-optimization/78024] Clear basic block flags before using BB_VISITED for OpenACC loops processing
Thomas Schwinge
thomas@codesourcery.com
Tue Jan 10 13:05:00 GMT 2017
Hi!
On Wed, 19 Oct 2016 12:28:39 +0200, I wrote:
> [PR tree-optimization/78024] Clear basic block flags before using BB_VISITED for OpenACC loops processing
To fix the same problem, committed the following to gcc-6-branch in
r244264:
commit 82f25a52380b366a99af8045fb615d79d4ff94b6
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Tue Jan 10 13:00:31 2017 +0000
[PR tree-optimization/78024] Clear basic block flags before using BB_VISITED for OpenACC loops processing
gcc/
PR tree-optimization/78024
* omp-low.c (oacc_loop_discovery): Call clear_bb_flags.
Backport from trunk r241334:
gcc/testsuite/
2016-10-19 Thomas Schwinge <thomas@codesourcery.com>
PR tree-optimization/78024
* gcc.dg/goacc/loop-processing-1.c: New file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gcc-6-branch@244264 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 3 +++
gcc/omp-low.c | 9 +++++----
gcc/testsuite/ChangeLog | 8 ++++++++
gcc/testsuite/gcc.dg/goacc/loop-processing-1.c | 18 ++++++++++++++++++
4 files changed, 34 insertions(+), 4 deletions(-)
diff --git gcc/ChangeLog gcc/ChangeLog
index 86b7f1b..5e7b2fc 100644
--- gcc/ChangeLog
+++ gcc/ChangeLog
@@ -1,5 +1,8 @@
2017-01-10 Thomas Schwinge <thomas@codesourcery.com>
+ PR tree-optimization/78024
+ * omp-low.c (oacc_loop_discovery): Call clear_bb_flags.
+
Backport trunk r239086:
2016-08-03 Nathan Sidwell <nathan@codesourcery.com>
diff --git gcc/omp-low.c gcc/omp-low.c
index 15ecb44..57a03df 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -19225,7 +19225,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));
@@ -19234,9 +19236,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. */
+ clear_bb_flags ();
return top;
}
diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog
index 391941f..df035a5 100644
--- gcc/testsuite/ChangeLog
+++ gcc/testsuite/ChangeLog
@@ -1,3 +1,11 @@
+2017-01-10 Thomas Schwinge <thomas@codesourcery.com>
+
+ Backport from trunk r241334:
+ 2016-10-19 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR tree-optimization/78024
+ * gcc.dg/goacc/loop-processing-1.c: New file.
+
2017-01-09 Andre Vieira <andre.simoesdiasvieira@arm.com>
Backport from mainline
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..2064bcd
--- /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 \\\(\[0-9\]+, 0, 1, 20\\\);.*Head-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(\[0-9\]+, 0, 1, 20\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(\[0-9\]+, \\\.data_dep\\\.\[0-9_\]+, 0\\\);.*Tail-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(\[0-9\]+, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(\[0-9\]+, \\\.data_dep\\\.\[0-9_\]+, 0\\\);.*Loop 6\\\(4\\\).*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(\[0-9\]+, 0, 1, 6\\\);.*Head-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(\[0-9\]+, 0, 1, 6\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(\[0-9\]+, \\\.data_dep\\\.\[0-9_\]+, 2\\\);.*Tail-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(\[0-9\]+, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(\[0-9\]+, \\\.data_dep\\\.\[0-9_\]+, 2\\\);" "oaccdevlow" } } */
..., and to gomp-4_0-branch in r244265:
commit 3b218c043a68883431fc5f3497d4e43137db8d8b
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Tue Jan 10 13:02:00 2017 +0000
[PR tree-optimization/78024] Clear basic block flags before using BB_VISITED for OpenACC loops processing
Backport from gcc-6-branch r244264:
gcc/
2017-01-10 Thomas Schwinge <thomas@codesourcery.com>
PR tree-optimization/78024
* omp-low.c (oacc_loop_discovery): Call clear_bb_flags.
Backport from trunk r241334:
gcc/testsuite/
2016-10-19 Thomas Schwinge <thomas@codesourcery.com>
PR tree-optimization/78024
* gcc.dg/goacc/loop-processing-1.c: New file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@244265 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 8 ++++++++
gcc/omp-low.c | 9 +++++----
gcc/testsuite/ChangeLog.gomp | 8 ++++++++
gcc/testsuite/gcc.dg/goacc/loop-processing-1.c | 18 ++++++++++++++++++
4 files changed, 39 insertions(+), 4 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index ecd45ec..c8a13d5 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,11 @@
+2017-01-10 Thomas Schwinge <thomas@codesourcery.com>
+
+ Backport from gcc-6-branch r244264:
+ 2017-01-10 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR tree-optimization/78024
+ * omp-low.c (oacc_loop_discovery): Call clear_bb_flags.
+
2017-01-10 Chung-Lin Tang <cltang@codesourcery.com>
* gimplify.c (gimplify_scan_omp_clauses): For dynamic array map kinds,
diff --git gcc/omp-low.c gcc/omp-low.c
index eee30bd..38c8f34 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -20359,7 +20359,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));
@@ -20368,9 +20370,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. */
+ clear_bb_flags ();
return top;
}
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 039a703..3e0e576 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,11 @@
+2017-01-10 Thomas Schwinge <thomas@codesourcery.com>
+
+ Backport from trunk r241334:
+ 2016-10-19 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR tree-optimization/78024
+ * gcc.dg/goacc/loop-processing-1.c: New file.
+
2016-12-06 Cesar Philippidis <cesar@codesourcery.com>
* c-c++-common/goacc/acc-data-chain.c: New test.
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..ac886c7
--- /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 24\\\(1\\\).*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, 0, 1, 36\\\);.*Head-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, 0, 1, 36\\\);.*\\\.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\\\(6\\\).*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, 0, 2, 6\\\);.*Head-0:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, 0, 2, 6\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_FORK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*Head-1:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_HEAD_MARK, \\\.data_dep\\\.\[0-9_\]+, 1\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_FORK, \\\.data_dep\\\.\[0-9_\]+, 2\\\);.*Tail-1:.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_TAIL_MARK, \\\.data_dep\\\.\[0-9_\]+, 2\\\);.*\\\.data_dep\\\.\[0-9_\]+ = UNIQUE \\\(OACC_JOIN, \\\.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_\]+, 1\\\);" "oaccdevlow" } } */
Grüße
Thomas
More information about the Gcc-patches
mailing list