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.
>
> On 2019-02-01T00:59:30+0100, I wrote:
>> 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.
>
> That's still the case... :-)
(The pass is still disabled by default, by the way.)
We've found that 'gcc/omp-oacc-kernels-decompose.cc' is currently not at
all considerate of 'GIMPLE_DEBUG' statements -- and it's not always
straight forward how to handle these (not rocket science either; but
needs proper understanding and testing).
Actually fixing it is a separate task, but it seems prudent to at least
catch it, and document via a few test cases. OK to push
"Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels' decomposition
[PR100400, PR103836, PR104061]", see attached?
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634
München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas
Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht
München, HRB 106955
>From 568808ef7ccc97ebeae90bc7cb1aba6bd7659b24 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Wed, 19 Jan 2022 14:04:42 +0100
Subject: [PATCH] Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels'
decomposition [PR100400, PR103836, PR104061]
Actually fixing it is a separate task, but it seems prudent to at least catch
it, and document via a few test cases.
gcc/
PR middle-end/100400
PR middle-end/103836
PR middle-end/104061
* omp-oacc-kernels-decompose.cc (decompose_kernels_region_body):
Catch 'GIMPLE_DEBUG'.
gcc/testsuite/
PR middle-end/100400
PR middle-end/103836
PR middle-end/104061
* c-c++-common/goacc/kernels-decompose-pr100400-1-1.c: New.
* c-c++-common/goacc/kernels-decompose-pr100400-1-2.c: New.
* c-c++-common/goacc/kernels-decompose-pr100400-1-3.c: New.
* c-c++-common/goacc/kernels-decompose-pr100400-1-4.c: New.
* c-c++-common/goacc/kernels-decompose-pr103836-1-1.c: New.
* c-c++-common/goacc/kernels-decompose-pr103836-1-2.c: New.
* c-c++-common/goacc/kernels-decompose-pr103836-1-3.c: New.
* c-c++-common/goacc/kernels-decompose-pr103836-1-4.c: New.
* c-c++-common/goacc/kernels-decompose-pr104061-1-1.c: New.
* c-c++-common/goacc/kernels-decompose-pr104061-1-2.c: New.
* c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: New.
* c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: New.
---
gcc/omp-oacc-kernels-decompose.cc | 10 +++++
.../goacc/kernels-decompose-pr100400-1-1.c | 33 ++++++++++++++
.../goacc/kernels-decompose-pr100400-1-2.c | 40 +++++++++++++++++
.../goacc/kernels-decompose-pr100400-1-3.c | 42 ++++++++++++++++++
.../goacc/kernels-decompose-pr100400-1-4.c | 40 +++++++++++++++++
.../goacc/kernels-decompose-pr103836-1-1.c | 26 +++++++++++
.../goacc/kernels-decompose-pr103836-1-2.c | 29 +++++++++++++
.../goacc/kernels-decompose-pr103836-1-3.c | 30 +++++++++++++
.../goacc/kernels-decompose-pr103836-1-4.c | 30 +++++++++++++
.../goacc/kernels-decompose-pr104061-1-1.c | 30 +++++++++++++
.../goacc/kernels-decompose-pr104061-1-2.c | 33 ++++++++++++++
.../goacc/kernels-decompose-pr104061-1-3.c | 43 +++++++++++++++++++
.../goacc/kernels-decompose-pr104061-1-4.c | 41 ++++++++++++++++++
13 files changed, 427 insertions(+)
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-2.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-3.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-4.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-2.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-3.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-4.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c
diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc
index 21872db3ed3..98eafdbe3a1 100644
--- a/gcc/omp-oacc-kernels-decompose.cc
+++ b/gcc/omp-oacc-kernels-decompose.cc
@@ -1255,6 +1255,16 @@ decompose_kernels_region_body (gimple *kernels_region, tree kernels_clauses)
gsi_next (&gsi_n);
gimple *stmt = gsi_stmt (gsi);
+ if (gimple_code (stmt) == GIMPLE_DEBUG)
+ {
+ if (flag_compare_debug_opt || flag_compare_debug)
+ /* Let the usual '-fcompare-debug' analysis bail out, as
+ necessary. */
+ ;
+ else
+ sorry_at (loc, "%qs not yet supported",
+ gimple_code_name[gimple_code (stmt)]);
+ }
gimple *omp_for = top_level_omp_for_in_stmt (stmt);
bool is_unconditional_oacc_for_loop = false;
if (omp_for != NULL)
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-1.c
new file mode 100644
index 00000000000..f63800514c4
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-1.c
@@ -0,0 +1,33 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-g0" } */
+/* { dg-additional-options "-O1" } */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+int *p;
+
+void
+foo (void)
+{
+#pragma acc kernels
+ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ {
+ int c;
+
+ /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+ p = &c;
+
+ /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+#pragma acc loop independent
+ /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } .-2 }
+ { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } .-3 } */
+ /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } .-4 } */
+ for (c = 0; c < 1; ++c)
+ ;
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-2.c
new file mode 100644
index 00000000000..1eee3b07a75
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-2.c
@@ -0,0 +1,40 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-fchecking" }
+ { dg-ice TODO { c++ } }
+ { dg-prune-output "during GIMPLE pass: omp_oacc_kernels_decompose" } */
+
+/* { dg-additional-options "-g" } */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+int *p;
+
+void
+foo (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail *-*-* } .+1 } */
+#pragma acc kernels
+ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { xfail *-*-* } .-1 } */
+ /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-2 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 }
+ { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */
+ int c;
+
+ /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { xfail *-*-* } .+1 } */
+ p = &c;
+
+ /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { xfail c++ } .+1 } */
+#pragma acc loop independent
+ /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-1 } */
+ /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { xfail *-*-* } .-2 }
+ { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { xfail *-*-* } .-3 } */
+ /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { xfail *-*-* } .-4 } */
+ for (c = 0; c < 1; ++c)
+ ;
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-3.c
new file mode 100644
index 00000000000..dce4e399fbe
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-3.c
@@ -0,0 +1,42 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-fchecking" }
+ { dg-ice TODO { c++ } }
+ { dg-prune-output "during GIMPLE pass: omp_oacc_kernels_decompose" } */
+
+/* { dg-additional-options "-fcompare-debug" } -- w/o debug compiled first.
+ { dg-bogus {error: during '-fcompare-debug' recompilation} TODO { xfail c++ } 0 }
+ { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail c++ } 0 } */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+int *p;
+
+void
+foo (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */
+#pragma acc kernels
+ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 }
+ { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */
+ int c;
+
+ /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+ p = &c;
+
+ /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+#pragma acc loop independent
+ /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } .-2 }
+ { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } .-3 } */
+ /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } .-4 } */
+ for (c = 0; c < 1; ++c)
+ ;
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-4.c
new file mode 100644
index 00000000000..7ca4440d075
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-4.c
@@ -0,0 +1,40 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-fchecking" }
+ { dg-ice TODO { c++ } }
+ { dg-prune-output "during GIMPLE pass: omp_oacc_kernels_decompose" } */
+
+/* { dg-additional-options "-g -fcompare-debug" } -- w/ debug compiled first. */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+int *p;
+
+void
+foo (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */
+#pragma acc kernels
+ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { xfail c++ } .-1 } */
+ /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-2 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 }
+ { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */
+ int c;
+
+ /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { xfail c++ } .+1 } */
+ p = &c;
+
+ /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { xfail c++ } .+1 } */
+#pragma acc loop independent
+ /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-1 } */
+ /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { xfail c++ } .-2 }
+ { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { xfail c++ } .-3 } */
+ /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { xfail c++ } .-4 } */
+ for (c = 0; c < 1; ++c)
+ ;
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-1.c
new file mode 100644
index 00000000000..46ca0c99d2f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-1.c
@@ -0,0 +1,26 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-g0" } */
+/* { dg-additional-options "-O1" } */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+extern int i;
+
+void
+f_acc_kernels (void)
+{
+#pragma acc kernels
+ /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ {
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-3 } */
+ for (i = 0; i < 2; ++i)
+ ;
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-2.c
new file mode 100644
index 00000000000..e0f24cee2db
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-2.c
@@ -0,0 +1,29 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-g" } */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+extern int i;
+
+void
+f_acc_kernels (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail c++ } .+1 } */
+#pragma acc kernels
+ /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-1 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } */
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-1 } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-2 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail c++ } .-3 } */
+ for (i = 0; i < 2; ++i)
+ ;
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-3.c
new file mode 100644
index 00000000000..cbf1b7c3e25
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-3.c
@@ -0,0 +1,30 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-fcompare-debug" } -- w/o debug compiled first.
+ { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail c++ } 0 } */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+extern int i;
+
+void
+f_acc_kernels (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } */
+#pragma acc kernels
+ /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } */
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-3 } */
+ for (i = 0; i < 2; ++i)
+ ;
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-4.c
new file mode 100644
index 00000000000..21bbe37723f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-4.c
@@ -0,0 +1,30 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-g -fcompare-debug" } -- w/ debug compiled first.
+ { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail c++ } 0 } */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+extern int i;
+
+void
+f_acc_kernels (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } */
+#pragma acc kernels
+ /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } */
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-3 } */
+ for (i = 0; i < 2; ++i)
+ ;
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-1.c
new file mode 100644
index 00000000000..a58fce33426
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-1.c
@@ -0,0 +1,30 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-g0" } */
+/* { dg-additional-options "-O1" } */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" }
+ Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
+ { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
+
+int arr_0;
+
+void
+foo (void)
+{
+#pragma acc kernels
+ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ {
+ int k;
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-3 } */
+ for (k = 0; k < 2; k++)
+ arr_0 += k;
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c
new file mode 100644
index 00000000000..d66dee6f8a7
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c
@@ -0,0 +1,33 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-g" } */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" }
+ Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
+ { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
+
+int arr_0;
+
+void
+foo (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail *-*-* } .+1 } */
+#pragma acc kernels
+ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-1 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 }
+ { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */
+ int k;
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-1 } */
+ /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-2 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail *-*-* } .-3 } */
+ for (k = 0; k < 2; k++)
+ arr_0 += k;
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c
new file mode 100644
index 00000000000..20c84e2f3db
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c
@@ -0,0 +1,43 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-fchecking" }
+ { dg-ice TODO }
+ { dg-prune-output {D\.[0-9]+ = arr_0\.0 \+ k;} }
+ { dg-prune-output {during GIMPLE pass: lower} } */
+
+/* { dg-additional-options "-fcompare-debug" } -- w/o debug compiled first.
+ { dg-bogus {error: during '-fcompare-debug' recompilation} TODO { xfail *-*-* } 0 }
+ { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail *-*-* } 0 } */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" }
+ Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
+ { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
+
+int arr_0;
+
+void
+foo (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */
+#pragma acc kernels
+ /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } .-1 } */
+ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 }
+ { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */
+ int k;
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ /* { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } .-3 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-4 } */
+ for (k = 0; k < 2; k++)
+ arr_0 += k;
+ /* { dg-bogus {error: invalid operands in binary operation} {w/ debug} { xfail *-*-* } .-1 } */
+ }
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c
new file mode 100644
index 00000000000..6b6effe1791
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c
@@ -0,0 +1,41 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-fchecking" }
+ { dg-ice TODO }
+ { dg-prune-output {D\.[0-9]+ = arr_0\.0 \+ k;} }
+ { dg-prune-output {during GIMPLE pass: lower} } */
+
+/* { dg-additional-options "-g -fcompare-debug" } -- w/ debug compiled first. */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" }
+ Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
+ { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
+
+int arr_0;
+
+void
+foo (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */
+#pragma acc kernels
+ /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } .-1 } */
+ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 }
+ { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */
+ int k;
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-1 } */
+ /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-2 } */
+ /* { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } .-3 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail *-*-* } .-4 } */
+ for (k = 0; k < 2; k++)
+ arr_0 += k;
+ /* { dg-bogus {error: invalid operands in binary operation} {w/ debug} { xfail *-*-* } .-1 } */
+ }
+}
--
2.25.1
