Hi!
On 2020-11-13T23:22:30+0100, I wrote:
> [...] I've pushed to master branch [...] commit
> e898ce7997733c29dcab9c3c62ca102c7f9fa6eb "Decompose OpenACC 'kernels'
> constructs into parts, a sequence of compute constructs", see attached.
>
>> There's more work to be done there, and we're aware of a number of TODO
>> items, but nevertheless: it's a good first step.
> --- /dev/null
> +++ b/gcc/omp-oacc-kernels-decompose.cc
> +/* Eliminate any binds directly inside BIND by adding their statements to
> + BIND (i.e., modifying it in place), excluding binds that hold only an
> + OMP_FOR loop and associated setup/cleanup code. Recurse into binds but
> + not other statements. Return a chain of the local variables of eliminated
> + binds, i.e., the local variables found in nested binds. If
> + INCLUDE_TOPLEVEL_VARS is true, this also includes the variables belonging
> + to BIND itself. */
> +
> +static tree
> +flatten_binds (gbind *bind, bool include_toplevel_vars = false)
> +{
> + tree vars = NULL, last_var = NULL;
> +
> + if (include_toplevel_vars)
> + {
> + vars = gimple_bind_vars (bind);
> + last_var = vars;
> + }
> +
> + gimple_seq new_body = NULL;
> + gimple_seq body_sequence = gimple_bind_body (bind);
> + gimple_stmt_iterator gsi, gsi_n;
> + for (gsi = gsi_start (body_sequence); !gsi_end_p (gsi); gsi = gsi_n)
> + {
> + /* Advance the iterator here because otherwise it would be invalidated
> + by moving statements below. */
> + gsi_n = gsi;
> + gsi_next (&gsi_n);
> +
> + gimple *stmt = gsi_stmt (gsi);
> + /* Flatten bind statements, except the ones that contain only an
> + OpenACC for loop. */
> + if (gimple_code (stmt) == GIMPLE_BIND
> + && !top_level_omp_for_in_stmt (stmt))
> + {
> + gbind *inner_bind = as_a <gbind *> (stmt);
> + /* Flatten recursively, and collect all variables. */
> + tree inner_vars = flatten_binds (inner_bind, true);
> + gimple_seq inner_sequence = gimple_bind_body (inner_bind);
> + gcc_assert (gimple_code (inner_sequence) != GIMPLE_BIND
> + || top_level_omp_for_in_stmt (inner_sequence));
First, this gives rise to the ICE documented in
'c-c++-common/goacc/kernels-decompose-ice-2.c': 'inner_sequence' is
'NULL' in this case (which is valid, meaning no statements in
'GIMPLE_BIND'), but 'gimple_code' then attempts to dereference 'NULL';
SIGSEGV.
Second, it seems strange to examine only the first statement of inner
'GIMPLE_BIND' (via 'inner_sequence' being a 'typedef gimple *gimple_seq'),
so I changed that to examine all statements contained therein, which I
suppose must've been the intention here. (This also conveniently fixes
the ICE mentioned above.)
I've pushed "In 'gcc/omp-oacc-kernels-decompose.cc:flatten_binds', don't
choke on empty GIMPLE sequence" to master branch in commit
4b5726fda653d11f882fb9a112e4cffa12f7ed61, see attached.
Grüße
Thomas
> + gimple_seq_add_seq (&new_body, inner_sequence);
> + /* Find the last variable; we will append others to it. */
> + while (last_var != NULL && TREE_CHAIN (last_var) != NULL)
> + last_var = TREE_CHAIN (last_var);
> + if (last_var != NULL)
> + {
> + TREE_CHAIN (last_var) = inner_vars;
> + last_var = inner_vars;
> + }
> + else
> + {
> + vars = inner_vars;
> + last_var = vars;
> + }
> + }
> + else
> + gimple_seq_add_stmt (&new_body, stmt);
> + }
> +
> + /* Put the possibly transformed body back into the bind. */
> + gimple_bind_set_body (bind, new_body);
> + return vars;
> +}
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-2.c
> @@ -0,0 +1,16 @@
> +/* Test OpenACC 'kernels' construct decomposition. */
> +
> +/* { dg-additional-options "-fopenacc-kernels=decompose" } */
> +/* { dg-ice "TODO" }
> + { dg-prune-output "during GIMPLE pass: omp_oacc_kernels_decompose" } */
> +
> +/* Reduced from 'kernels-decompose-ice-1.c'. */
> +
> +int
> +main ()
> +{
> +#pragma acc kernels
> + {
> + int i;
> + }
> +}
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander
Walter
>From 4b5726fda653d11f882fb9a112e4cffa12f7ed61 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Fri, 27 Nov 2020 11:54:50 +0100
Subject: [PATCH] In 'gcc/omp-oacc-kernels-decompose.cc:flatten_binds', don't
choke on empty GIMPLE sequence
Also, instead of just examining the first statement of inner 'GIMPLE_BIND' (via
'inner_sequence' being a 'typedef gimple *gimple_seq'), in fact examine all
statements contained therein, which I suppose must've been the intention here.
This "fixes" the testcase 'c-c++-common/goacc/kernels-decompose-ice-2.c' (which
now runs into the same ICE as 'c-c++-common/goacc/kernels-decompose-ice-1.c',
etc.).
gcc/
* omp-oacc-kernels-decompose.cc (flatten_binds): Don't choke on
empty GIMPLE sequence, and examine all statements contained in
inner 'GIMPLE_BIND'.
gcc/testsuite/
* c-c++-common/goacc/kernels-decompose-ice-1.c: Adjust.
* c-c++-common/goacc/kernels-decompose-ice-2.c: Likewise.
---
gcc/omp-oacc-kernels-decompose.cc | 13 +++++++++++--
.../c-c++-common/goacc/kernels-decompose-ice-1.c | 1 +
.../c-c++-common/goacc/kernels-decompose-ice-2.c | 2 +-
3 files changed, 13 insertions(+), 3 deletions(-)
diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc
index baad1b9a348..c46168e063a 100644
--- a/gcc/omp-oacc-kernels-decompose.cc
+++ b/gcc/omp-oacc-kernels-decompose.cc
@@ -740,8 +740,17 @@ flatten_binds (gbind *bind, bool include_toplevel_vars = false)
/* Flatten recursively, and collect all variables. */
tree inner_vars = flatten_binds (inner_bind, true);
gimple_seq inner_sequence = gimple_bind_body (inner_bind);
- gcc_assert (gimple_code (inner_sequence) != GIMPLE_BIND
- || top_level_omp_for_in_stmt (inner_sequence));
+ if (flag_checking)
+ {
+ for (gimple_stmt_iterator inner_gsi = gsi_start (inner_sequence);
+ !gsi_end_p (inner_gsi);
+ gsi_next (&inner_gsi))
+ {
+ gimple *inner_stmt = gsi_stmt (inner_gsi);
+ gcc_assert (gimple_code (inner_stmt) != GIMPLE_BIND
+ || top_level_omp_for_in_stmt (inner_stmt));
+ }
+ }
gimple_seq_add_seq (&new_body, inner_sequence);
/* Find the last variable; we will append others to it. */
while (last_var != NULL && TREE_CHAIN (last_var) != NULL)
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-1.c
index 9e27d1fb9b5..82e7bd1495b 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-1.c
@@ -7,6 +7,7 @@
/* Reduced from 'kernels-decompose-2.c'.
(Hopefully) similar instances:
+ - 'kernels-decompose-ice-2.c'
- 'libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c'
- 'libgomp.oacc-c-c++-common/kernels-decompose-1.c'
*/
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-2.c
index 839e6803851..569f87a59c9 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-ice-2.c
@@ -2,7 +2,7 @@
/* { dg-additional-options "-fopenacc-kernels=decompose" } */
/* { dg-ice "TODO" }
- { dg-prune-output "during GIMPLE pass: omp_oacc_kernels_decompose" } */
+ { dg-prune-output "during GIMPLE pass: omplower" } */
/* Reduced from 'kernels-decompose-ice-1.c'. */
--
2.17.1
