omp-oacc-kernels-decompose.cc: fix -fcompare-debug with GIMPLE_DEBUG

GIMPLE_DEBUG were put in a parallel region of its own, which is not
only pointless but also breaks -fcompare-debug. With this commit,
they are handled like simple assignments: those placed are places
into the same body as the loop such that only one parallel region
remains as without debugging. This fixes the existing testcase
libgomp.oacc-c-c++-common/kernels-loop-g.c.

Note: GIMPLE_DEBUG are only accepted with -fcompare-debug; if they
appear otherwise, decompose_kernels_region_body rejects them with
a sorry (unchanged).

Also note that there are still many xfailed tests in the
c-c++-common/goacc/kernels-decompose-pr* testcases that were added
in mainline commit c14ea6a72fb1ae66e3d32ac8329558497c6e4403.

gcc/ChangeLog
	* omp-oacc-kernels-decompose.cc (top_level_omp_for_in_stmt,
	decompose_kernels_region_body): Handle GIMPLE_DEBUG like
	simple assignment.

gcc/testsuite/ChangeLog
	* c-c++-common/goacc/kernels-decompose-pr103836-1-2.c: Adjust xfails.
	* c-c++-common/goacc/kernels-decompose-pr103836-1-3.c: Likewise.
	* c-c++-common/goacc/kernels-decompose-pr103836-1-4.c: Likewise.
	* c-c++-common/goacc/kernels-decompose-pr104061-1-2.c: Likewise.
	* c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: Likewise.
	* c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: Likewise.

Co-Authored-By: Sandra Loosemore <sloosemore@baylibre.com>

7 files changed, 29 insertions, 28 deletions

diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc
index fc3a3b3..5d7289a 100644
--- a/gcc/omp-oacc-kernels-decompose.cc
+++ b/gcc/omp-oacc-kernels-decompose.cc
@@ -120,7 +120,8 @@ top_level_omp_for_in_stmt (gimple *stmt)
 	  for (gsi = gsi_start (body); !gsi_end_p (gsi); gsi_next (&gsi))
 	    {
 	      gimple *body_stmt = gsi_stmt (gsi);
-	      if (gimple_code (body_stmt) == GIMPLE_ASSIGN)
+	      if (gimple_code (body_stmt) == GIMPLE_ASSIGN
+		  || gimple_code (body_stmt) == GIMPLE_DEBUG)
 		continue;
 	      else if (gimple_code (body_stmt) == GIMPLE_OMP_FOR
 		       && gsi_one_before_end_p (gsi))
@@ -1363,7 +1364,7 @@ decompose_kernels_region_body (gimple *kernels_region, tree kernels_clauses)
 	    = (gimple_code (stmt) == GIMPLE_ASSIGN
 	       && TREE_CODE (gimple_assign_lhs (stmt)) == VAR_DECL
 	       && DECL_ARTIFICIAL (gimple_assign_lhs (stmt)));
-	  if (!is_simple_assignment)
+	  if (!is_simple_assignment && gimple_code (stmt) != GIMPLE_DEBUG)
 	    only_simple_assignments = false;
 	}
     }
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
index 83690b6..a3afb79 100644
--- 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
@@ -16,7 +16,7 @@ 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} {} { 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} { target *-*-* } .-1 } */
 
     /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
 #pragma acc loop /* { dg-line l_loop_i1 } */
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
index 35892a0..8cbf69f 100644
--- 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
@@ -1,7 +1,7 @@
 /* { 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-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} "" { target *-*-* } 0 } */
 /* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's.  */
 
 /* { dg-additional-options "-fopt-info-all-omp" } */
@@ -17,7 +17,7 @@ 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-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} { target *-*-* } .-1 } */
 
     /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
 #pragma acc loop /* { dg-line l_loop_i1 } */
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
index 549ad5d..507e73c 100644
--- 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
@@ -1,7 +1,7 @@
 /* { 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-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} "" { target *-*-* } 0 } */
 /* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's.  */
 
 /* { dg-additional-options "-fopt-info-all-omp" } */
@@ -17,7 +17,7 @@ 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-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} { target *-*-* } .-1 } */
 
     /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
 #pragma acc loop /* { dg-line l_loop_i1 } */
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
index 4d7cbb0..6865a5c 100644
--- 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
@@ -14,23 +14,23 @@ int arr_0;
 void
 foo (void)
 {
-  /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail *-*-* } .+1 } */
+  /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail c++ } .+1 } */
 #pragma acc kernels /* { dg-line l_compute1 } */
   /* { dg-note {OpenACC 'kernels' decomposition: variable 'arr_0' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 }
      { dg-note {variable 'arr_0' made addressable} {} { target *-*-* } l_compute1 } */
-  /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */
-  /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { xfail *-*-* } 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: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { target *-*-* } l_compute1 } */
+  /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { target *-*-* } l_compute1 } */
+  /* { dg-note {variable 'arr_0\.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 } */
+    /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { target c++ } .-1 }
+       { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { target 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 } */
+    /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++} l_loop_k1 } */
+    /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } l_loop_k1 } */
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail c++ } 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
index 70c2ac5..197cee3 100644
--- 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
@@ -1,7 +1,7 @@
 /* { 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 *-*-* } 0 } */
+   { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} "" { target *-*-* } 0 } */
 /* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's.  */
 
 /* { dg-additional-options "-fopt-info-all-omp" } */
@@ -19,13 +19,13 @@ foo (void)
 #pragma acc kernels /* { dg-line l_compute1 } */
   /* { dg-note {OpenACC 'kernels' decomposition: variable 'arr_0' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 }
      { dg-note {variable 'arr_0' made addressable} {} { target *-*-* } l_compute1 } */
-  /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */
-  /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */
-  /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_compute1 } */
+  /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { target *-*-* } l_compute1 } */
+  /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { target *-*-* } l_compute1 } */
+  /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { target *-*-* } 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 } */
+    /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { target c++ } .-1 }
+       { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { target c } .+1 } */
     int k;
 
     /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+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
index d1cc1a9..f82c1c7 100644
--- 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
@@ -1,7 +1,7 @@
 /* { 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 *-*-* } 0 } */
+   { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} "" { target *-*-* } 0 } */
 /* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's.  */
 
 /* { dg-additional-options "-fopt-info-all-omp" } */
@@ -19,13 +19,13 @@ foo (void)
 #pragma acc kernels /* { dg-line l_compute1 } */
   /* { dg-note {OpenACC 'kernels' decomposition: variable 'arr_0' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 }
      { dg-note {variable 'arr_0' made addressable} {} { target *-*-* } l_compute1 } */
-  /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */
-  /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */
-  /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_compute1 } */
+  /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { target *-*-* } l_compute1 } */
+  /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { target *-*-* } l_compute1 } */
+  /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { target *-*-* } 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 } */
+    /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { target c++ } .-1 }
+       { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { target c } .+1 } */
     int k;
 
     /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */