From 568808ef7ccc97ebeae90bc7cb1aba6bd7659b24 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@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

