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.

13 files changed, 427 insertions, 0 deletions

diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc
index 21872db..98eafdb 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 0000000..b1b094f
--- /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-line l_compute1 } */
+  /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */
+  /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */
+  {
+    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-line l_loop_c1 } */
+    /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_c1 } */
+    /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_c1 }
+       { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l_loop_c1 } */
+    /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l_loop_c1 } */
+    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 0000000..0bc4844
--- /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-line l_compute1 } */
+  /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { xfail *-*-* } l_compute1 } */
+  /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_compute1 } */
+  {
+    /* { 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-line l_loop_c1 } */
+    /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_loop_c1 } */
+    /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { xfail *-*-* } l_loop_c1 }
+       { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { xfail *-*-* } l_loop_c1 } */
+    /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { xfail *-*-* } l_loop_c1 } */
+    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 0000000..6c2cbf8
--- /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-line l_compute1 } */
+  /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */
+  /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */
+  {
+    /* { 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-line l_loop_c1 } */
+    /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_c1 } */
+    /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_c1 }
+       { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l_loop_c1 } */
+    /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l_loop_c1 } */
+    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 0000000..fc66a3e
--- /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-line l_compute1 } */
+  /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { xfail c++ } l_compute1 } */
+  /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } l_compute1 } */
+  {
+    /* { 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-line l_loop_c1 } */
+    /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } l_loop_c1 } */
+    /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { xfail c++ } l_loop_c1 }
+       { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { xfail c++ } l_loop_c1 } */
+    /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { xfail c++ } l_loop_c1 } */
+    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 0000000..253e4d6
--- /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-line l_compute1 } */
+  /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */
+  {
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop /* { dg-line l_loop_i1 } */
+    /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i1 } */
+    /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i1 } */
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i1 } */
+    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 0000000..83690b6
--- /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-line l_compute1 } */
+  /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } l_compute1 } */
+  {
+    /* { 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-line l_loop_i1 } */
+    /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } l_loop_i1 } */
+    /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } l_loop_i1 } */
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail c++ } l_loop_i1 } */
+    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 0000000..35892a0
--- /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-line l_compute1 } */
+  /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */
+  {
+    /* { 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-line l_loop_i1 } */
+    /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i1 } */
+    /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i1 } */
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i1 } */
+    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 0000000..549ad5d
--- /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-line l_compute1 } */
+  /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */
+  {
+    /* { 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-line l_loop_i1 } */
+    /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i1 } */
+    /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i1 } */
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i1 } */
+    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 0000000..35ae81f
--- /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-line l_compute1 } */
+  /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */
+  {
+    int k;
+
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop /* { dg-line l_loop_k1 } */
+    /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k1 } */
+    /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k1 } */
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_k1 } */
+    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 0000000..ea58ea4
--- /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-line l_compute1 } */
+  /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_compute1 } */
+  {
+    /* { 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-line l_loop_k1 } */
+    /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_loop_k1 } */
+    /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_loop_k1 } */
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail *-*-* } l_loop_k1 } */
+    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 0000000..ee3fcc7
--- /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-line l_compute1 } */
+  /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_compute1 } */
+  /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */
+  {
+    /* { 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-line l_loop_k1 } */
+    /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k1 } */
+    /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k1 } */
+    /* { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_loop_k1 } */
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_k1 } */
+    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 0000000..c140e61
--- /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-line l_compute1 } */
+  /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_compute1 } */
+  /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */
+  {
+    /* { 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-line l_loop_k1 } */
+    /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_loop_k1 } */
+    /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_loop_k1 } */
+    /* { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_loop_k1 } */
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail *-*-* } l_loop_k1 } */
+    for (k = 0; k < 2; k++)
+      arr_0 += k;
+      /* { dg-bogus {error: invalid operands in binary operation} {w/ debug} { xfail *-*-* } .-1 } */
+  }
+}