Test 'libgomp.oacc-*/kernels-private-vars-*' with '--param=openacc-kernels=decompose' [PR104784]

Before recent commit 8935589b496f755e08cadf26d8ceddf0dd6e0968
"OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs
[PR100280, PR104132, PR104133]", 'libgomp.oacc-c' testing already worked fine,
but 'libgomp.oacc-c++' testing ICEed.  Via the commit mentioned, the C++
testing ICEs are now resolved, but the underlying issue remains to be looked
into: PR104784 "OpenACC 'kernels' decomposition: C vs. C++ differences".

	PR middle-end/104784
	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-1.c:
	Test with '--param=openacc-kernels=decompose'.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-3.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-4.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-5.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-1.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-3.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-4.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-5.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-6.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-1.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-1.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-3.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-4.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-5.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-6.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-7.c:
	Likewise.
	* testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-1.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-2.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-3.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-6.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-1.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-2.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-1.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-2.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-3.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-4.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-5.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-6.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-7.f90:
	Likewise.

33 files changed, 361 insertions, 99 deletions

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-1.c
index acbeb65..43bfaf3 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-1.c
@@ -1,11 +1,15 @@
 /* Test of worker-private variables declared in a local scope, broadcasting
    to vector-partitioned mode.  Back-to-back worker loops.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -26,16 +30,24 @@ main (int argc, char* argv[])
     arr[i] = i;
 
   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
+     decides that a data region is needed for 'j', and subsequently requests it
+     to be made addressable.
+     { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
   {
     int j;
 
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     for (i = 0; i < 32; i++)
       {
 	#pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	for (j = 0; j < 32; j++)
@@ -50,7 +62,8 @@ main (int argc, char* argv[])
 	  }
 
 	#pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	for (j = 0; j < 32; j++)
@@ -64,8 +77,8 @@ main (int argc, char* argv[])
 	      arr[i * 1024 + j * 32 + k] += x * k;
 	  }
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32; i++)
     for (int j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-2.c
index 2558a68..c40c2ab 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-2.c
@@ -1,11 +1,15 @@
 /* Test of worker-private variables declared in a local scope, broadcasting
    to vector-partitioned mode.  Successive vector loops.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -26,16 +30,24 @@ main (int argc, char* argv[])
     arr[i] = i;
 
   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
+     decides that a data region is needed for 'j', and subsequently requests it
+     to be made addressable.
+     { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
   {
     int j;
 
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     for (i = 0; i < 32; i++)
       {
         #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	for (j = 0; j < 32; j++)
@@ -56,8 +68,8 @@ main (int argc, char* argv[])
 	      arr[i * 1024 + j * 32 + k] += x * k;
 	  }
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32; i++)
     for (int j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-3.c
index b2a208c..bd04dcc 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-3.c
@@ -1,11 +1,15 @@
 /* Test of worker-private variables declared in a local scope, broadcasting
    to vector-partitioned mode.  Aggregate worker variable.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -31,16 +35,24 @@ main (int argc, char* argv[])
     arr[i] = i;
 
   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
+     decides that a data region is needed for 'j', and subsequently requests it
+     to be made addressable.
+     { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
   {
     int j;
 
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     for (i = 0; i < 32; i++)
       {
         #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'pt' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	for (j = 0; j < 32; j++)
@@ -62,8 +74,8 @@ main (int argc, char* argv[])
 	      arr[i * 1024 + j * 32 + k] += pt.y * k;
 	  }
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32; i++)
     for (int j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-4.c
index 46c3956..4303ab8 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-4.c
@@ -1,11 +1,15 @@
 /* Test of worker-private variables declared in a local scope, broadcasting
    to vector-partitioned mode.  Addressable worker variable.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -31,16 +35,24 @@ main (int argc, char* argv[])
     arr[i] = i;
 
   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
+     decides that a data region is needed for 'j', and subsequently requests it
+     to be made addressable.
+     { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
   {
     int j;
 
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     for (i = 0; i < 32; i++)
       {
         #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'pt' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'ptp' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
@@ -66,8 +78,8 @@ main (int argc, char* argv[])
 	      arr[i * 1024 + j * 32 + k] += pt.y * k;
 	  }
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32; i++)
     for (int j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-5.c
index 4b5a15e..8d0e846 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-5.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-5.c
@@ -1,11 +1,15 @@
 /* Test of worker-private variables declared in a local scope, broadcasting
    to vector-partitioned mode.  Array worker variable.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -26,16 +30,24 @@ main (int argc, char* argv[])
     arr[i] = i;
 
   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
+     decides that a data region is needed for 'j', and subsequently requests it
+     to be made addressable.
+     { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
   {
     int j;
 
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     for (i = 0; i < 32; i++)
       {
         #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'pt' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	for (j = 0; j < 32; j++)
@@ -58,8 +70,8 @@ main (int argc, char* argv[])
 	      arr[i * 1024 + j * 32 + k] += pt[1] * k;
 	  }
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32; i++)
     for (int j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-1.c
index 4a82494..470d5e6 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-1.c
@@ -1,10 +1,14 @@
 /* Test of gang-private variables declared on loop directive.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -24,8 +28,9 @@ main (int argc, char* argv[])
   for (i = 0; i < 32; i++)
     arr[i] = i;
 
-  #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
+  #pragma acc kernels copy(arr)
   {
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) private(x) /* { dg-line l_loop_i[incr c_loop_i] } */
     /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
@@ -34,8 +39,8 @@ main (int argc, char* argv[])
 	x = i * 2;
 	arr[i] += x;
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32; i++)
     assert (arr[i] == i * 3);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-2.c
index 039053f..34513e3 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-2.c
@@ -1,11 +1,15 @@
 /* Test of gang-private variables declared on loop directive, with broadcasting
    to partitioned workers.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -25,8 +29,9 @@ main (int argc, char* argv[])
   for (i = 0; i < 32 * 32; i++)
     arr[i] = i;
 
-  #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
+  #pragma acc kernels copy(arr)
   {
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) private(x) /* { dg-line l_loop_i[incr c_loop_i] } */
     /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
@@ -40,8 +45,8 @@ main (int argc, char* argv[])
 	for (int j = 0; j < 32; j++)
 	  arr[i * 32 + j] += x;
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32 * 32; i++)
     assert (arr[i] == i + (i / 32) * 2);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-3.c
index 2b89659..ef445c8 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-3.c
@@ -1,11 +1,15 @@
 /* Test of gang-private variables declared on loop directive, with broadcasting
    to partitioned vectors.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -25,8 +29,9 @@ main (int argc, char* argv[])
   for (i = 0; i < 32 * 32; i++)
     arr[i] = i;
 
-  #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
+  #pragma acc kernels copy(arr)
   {
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) private(x) /* { dg-line l_loop_i[incr c_loop_i] } */
     /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
@@ -40,8 +45,8 @@ main (int argc, char* argv[])
 	for (int j = 0; j < 32; j++)
 	  arr[i * 32 + j] += x;
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32 * 32; i++)
     assert (arr[i] == i + (i / 32) * 2);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-4.c
index 7076070..0658251 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-4.c
@@ -1,11 +1,15 @@
 /* Test of gang-private addressable variable declared on loop directive, with
    broadcasting to partitioned workers.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -26,7 +30,9 @@ main (int argc, char* argv[])
     arr[i] = i;
 
   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
+  /* { dg-note {variable 'x\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
   {
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) private(x) /* { dg-line l_loop_i[incr c_loop_i] } */
     /* { dg-note {variable 'x' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
@@ -45,8 +51,8 @@ main (int argc, char* argv[])
 
 	(*p)--;
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32 * 32; i++)
     assert (arr[i] == i + (i / 32) * 2);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-5.c
index edf0e24..a511866 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-5.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-5.c
@@ -1,11 +1,15 @@
 /* Test of gang-private array variable declared on loop directive, with
    broadcasting to partitioned workers.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -25,8 +29,9 @@ main (int argc, char* argv[])
   for (i = 0; i < 32 * 32; i++)
     arr[i] = i;
 
-  #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
+  #pragma acc kernels copy(arr)
   {
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) private(x) /* { dg-line l_loop_i[incr c_loop_i] } */
     /* { dg-note {variable 'x' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
@@ -41,8 +46,8 @@ main (int argc, char* argv[])
 	for (int j = 0; j < 32; j++)
 	  arr[i * 32 + j] += x[j % 8];
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32 * 32; i++)
     assert (arr[i] == i + (i % 8) * 2);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-6.c
index a2df33b..119bda3 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-6.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-gang-6.c
@@ -1,11 +1,15 @@
 /* Test of gang-private aggregate variable declared on loop directive, with
    broadcasting to partitioned workers.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -31,8 +35,9 @@ main (int argc, char* argv[])
   for (i = 0; i < 32 * 32; i++)
     arr[i] = i;
 
-  #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
+  #pragma acc kernels copy(arr)
   {
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang private(pt) /* { dg-line l_loop_i[incr c_loop_i] } */
     /* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
@@ -49,8 +54,8 @@ main (int argc, char* argv[])
 	for (int j = 0; j < 32; j++)
 	  arr[i * 32 + j] += pt.x + pt.y + pt.z + pt.attr[5];
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32 * 32; i++)
     assert (arr[i] == i + (i / 32) * 13);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-1.c
index 51c1de5..5a70bb8 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-1.c
@@ -1,10 +1,14 @@
 /* Test of vector-private variables declared on loop directive.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -25,16 +29,24 @@ main (int argc, char* argv[])
     arr[i] = i;
 
   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
+     decides that a data region is needed for 'j', and subsequently requests it
+     to be made addressable.
+     { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
   {
     int j;
 
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     for (i = 0; i < 32; i++)
       {
         #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	for (j = 0; j < 32; j++)
 	  {
@@ -59,8 +71,8 @@ main (int argc, char* argv[])
 	      }
 	  }
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32; i++)
     for (int j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-2.c
index cb90eaa..f5bccab 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-2.c
@@ -1,10 +1,14 @@
 /* Test of vector-private variables declared on loop directive. Array type.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -25,16 +29,24 @@ main (int argc, char* argv[])
     arr[i] = i;
 
   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
+     decides that a data region is needed for 'j', and subsequently requests it
+     to be made addressable.
+     { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
   {
     int j;
 
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c} l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     for (i = 0; i < 32; i++)
       {
         #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	for (j = 0; j < 32; j++)
 	  {
@@ -52,8 +64,8 @@ main (int argc, char* argv[])
 	      }
 	  }
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32; i++)
     for (int j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-1.c
index 54e1c93..9b2c49a 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-1.c
@@ -1,10 +1,14 @@
 /* Test of worker-private variables declared on a loop directive.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -24,12 +28,13 @@ main (int argc, char* argv[])
   for (i = 0; i < 32 * 32; i++)
     arr[i] = i;
 
-  #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
+  #pragma acc kernels copy(arr)
   {
     int j;
 
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     for (i = 0; i < 32; i++)
       {
@@ -45,8 +50,8 @@ main (int argc, char* argv[])
 	    arr[i * 32 + j] += x;
 	  }
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32 * 32; i++)
     assert (arr[i] == i + ((i / 32) ^ (i % 32) * 3));
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-2.c
index 80ac990..40baae3 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-2.c
@@ -1,11 +1,15 @@
 /* Test of worker-private variables declared on a loop directive, broadcasting
    to vector-partitioned mode.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -26,17 +30,25 @@ main (int argc, char* argv[])
     arr[i] = i;
 
   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
+     decides that a data region is needed for 'j', and subsequently requests it
+     to be made addressable.
+     { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
   {
     int j;
 
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     for (i = 0; i < 32; i++)
       {
         #pragma acc loop worker(num:32) private(x) /* { dg-line l_loop_j[incr c_loop_j] } */
 	/* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	for (j = 0; j < 32; j++)
 	  {
@@ -49,8 +61,8 @@ main (int argc, char* argv[])
 	      arr[i * 1024 + j * 32 + k] += x * k;
 	  }
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32; i++)
     for (int j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-3.c
index a05ac60..c8b089c 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-3.c
@@ -1,11 +1,15 @@
 /* Test of worker-private variables declared on a loop directive, broadcasting
    to vector-partitioned mode.  Back-to-back worker loops.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -26,17 +30,25 @@ main (int argc, char* argv[])
     arr[i] = i;
 
   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
+     decides that a data region is needed for 'j', and subsequently requests it
+     to be made addressable.
+     { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
   {
     int j;
 
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     for (i = 0; i < 32; i++)
       {
         #pragma acc loop worker(num:32) private(x) /* { dg-line l_loop_j[incr c_loop_j] } */
 	/* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	for (j = 0; j < 32; j++)
 	  {
@@ -51,7 +63,8 @@ main (int argc, char* argv[])
 
 	#pragma acc loop worker(num:32) private(x) /* { dg-line l_loop_j[incr c_loop_j] } */
 	/* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	for (j = 0; j < 32; j++)
 	  {
@@ -64,8 +77,8 @@ main (int argc, char* argv[])
 	      arr[i * 1024 + j * 32 + k] += x * k;
 	  }
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32; i++)
     for (int j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-4.c
index d46bb94..c1819d2 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-4.c
@@ -1,11 +1,15 @@
 /* Test of worker-private variables declared on a loop directive, broadcasting
    to vector-partitioned mode.  Successive vector loops.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -26,17 +30,25 @@ main (int argc, char* argv[])
     arr[i] = i;
 
   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
+     decides that a data region is needed for 'j', and subsequently requests it
+     to be made addressable.
+     { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
   {
     int j;
 
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     for (i = 0; i < 32; i++)
       {
         #pragma acc loop worker(num:32) private(x) /* { dg-line l_loop_j[incr c_loop_j] } */
 	/* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	for (j = 0; j < 32; j++)
 	  {
@@ -56,8 +68,8 @@ main (int argc, char* argv[])
 	      arr[i * 1024 + j * 32 + k] += x * k;
 	  }
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32; i++)
     for (int j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-5.c
index 644c617..90955aa 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-5.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-5.c
@@ -1,11 +1,15 @@
 /* Test of worker-private variables declared on a loop directive, broadcasting
    to vector-partitioned mode.  Addressable worker variable.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -26,17 +30,26 @@ main (int argc, char* argv[])
     arr[i] = i;
 
   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
+     decides that a data region is needed for 'j', and subsequently requests it
+     to be made addressable.
+     { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
+  /* { dg-note {variable 'x\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
   {
     int j;
 
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     for (i = 0; i < 32; i++)
       {
         #pragma acc loop worker(num:32) private(x) /* { dg-line l_loop_j[incr c_loop_j] } */
 	/* { dg-note {variable 'x' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_j$c_loop_j } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'p' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	for (j = 0; j < 32; j++)
@@ -59,8 +72,8 @@ main (int argc, char* argv[])
 	      arr[i * 1024 + j * 32 + k] += x * k;
 	  }
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32; i++)
     for (int j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-6.c
index 182a12a..f093cfe 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-6.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-6.c
@@ -1,11 +1,15 @@
 /* Test of worker-private variables declared on a loop directive, broadcasting
    to vector-partitioned mode.  Aggregate worker variable.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -32,17 +36,25 @@ main (int argc, char* argv[])
     arr[i] = i;
 
   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
+     decides that a data region is needed for 'j', and subsequently requests it
+     to be made addressable.
+     { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
   {
     int j;
 
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     for (i = 0; i < 32; i++)
       {
         #pragma acc loop worker(num:32) private(pt) /* { dg-line l_loop_j[incr c_loop_j] } */
 	/* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	for (j = 0; j < 32; j++)
 	  {
@@ -62,8 +74,8 @@ main (int argc, char* argv[])
 	      arr[i * 1024 + j * 32 + k] += pt.y * k;
 	  }
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32; i++)
     for (int j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-7.c
index bdfbb59..906119c 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-7.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-7.c
@@ -1,11 +1,15 @@
 /* Test of worker-private variables declared on loop directive, broadcasting
    to vector-partitioned mode.  Array worker variable.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-omp-all" } */
 
 /* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
+   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+   Prune a few: uninteresting:
+   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
 
 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
    passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -29,18 +33,26 @@ main (int argc, char* argv[])
   /* "pt" is treated as "present_or_copy" on the kernels directive because it
      is an array variable.  */
   #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
+     decides that a data region is needed for 'j', and subsequently requests it
+     to be made addressable.
+     { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
+     { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
   {
     int j;
 
+    /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
     #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
+    /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
     /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
     for (i = 0; i < 32; i++)
       {
         /* But here, it is made private per-worker.  */
         #pragma acc loop worker(num:32) private(pt) /* { dg-line l_loop_j[incr c_loop_j] } */
 	/* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
-	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
+	/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_j$c_loop_j }
+	   { dg-note {variable 'j' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target c++ } l_loop_j$c_loop_j } */
 	/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
 	for (j = 0; j < 32; j++)
 	  {
@@ -61,8 +73,8 @@ main (int argc, char* argv[])
 	      arr[i * 1024 + j * 32 + k] += pt[1] * k;
 	  }
       }
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
   }
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
 
   for (i = 0; i < 32; i++)
     for (int j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-1.f90
index 09ab395..b4f44d6 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-1.f90
@@ -2,11 +2,15 @@
 
 ! { dg-do run }
 
+! { dg-additional-options "--param=openacc-kernels=decompose" }
+
 ! { dg-additional-options "-fopt-info-omp-all" }
 ! { dg-additional-options "-foffload=-fopt-info-omp-all" }
 
 ! { dg-additional-options "--param=openacc-privatization=noisy" }
 ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+! Prune a few: uninteresting:
+! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
 
 ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
 ! passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -23,7 +27,8 @@ program main
      arr(i) = i
   end do
 
-  !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] }
+  !$acc kernels copy(arr)
+  ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 }
   !$acc loop gang(num:32) private(x) ! { dg-line l_loop_i[incr c_loop_i] }
   ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
   ! { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
@@ -31,8 +36,8 @@ program main
      x = i * 2;
      arr(i) = arr(i) + x;
   end do
+  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i }
   !$acc end kernels
-  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute }
 
   do i = 1, 32
      if (arr(i) .ne. i * 3) stop 1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-2.f90
index bec1069..f4194c9 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-2.f90
@@ -3,11 +3,15 @@
 
 ! { dg-do run }
 
+! { dg-additional-options "--param=openacc-kernels=decompose" }
+
 ! { dg-additional-options "-fopt-info-omp-all" }
 ! { dg-additional-options "-foffload=-fopt-info-omp-all" }
 
 ! { dg-additional-options "--param=openacc-privatization=noisy" }
 ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+! Prune a few: uninteresting:
+! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
 
 ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
 ! passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -24,7 +28,8 @@ program main
      arr(i) = i
   end do
 
-  !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] }
+  !$acc kernels copy(arr)
+  ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 }
   !$acc loop gang(num:32) private(x) ! { dg-line l_loop_i[incr c_loop_i] }
   ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
   ! { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
@@ -37,8 +42,8 @@ program main
         arr(i * 32 + j) = arr(i * 32 + j) + x;
      end do
   end do
+  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i }
   !$acc end kernels
-  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute }
 
   do i = 0, 32 * 32 - 1
      if (arr(i) .ne. i + (i / 32) * 2) stop 1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-3.f90
index 9fde012..da85ac4 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-3.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-3.f90
@@ -3,11 +3,15 @@
 
 ! { dg-do run }
 
+! { dg-additional-options "--param=openacc-kernels=decompose" }
+
 ! { dg-additional-options "-fopt-info-omp-all" }
 ! { dg-additional-options "-foffload=-fopt-info-omp-all" }
 
 ! { dg-additional-options "--param=openacc-privatization=noisy" }
 ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+! Prune a few: uninteresting:
+! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
 
 ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
 ! passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -24,7 +28,8 @@ program main
      arr(i) = i
   end do
 
-  !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] }
+  !$acc kernels copy(arr)
+  ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 }
   !$acc loop gang(num:32) private(x) ! { dg-line l_loop_i[incr c_loop_i] }
   ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
   ! { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
@@ -37,8 +42,8 @@ program main
         arr(i * 32 + j) = arr(i * 32 + j) + x;
      end do
   end do
+  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i }
   !$acc end kernels
-  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute }
 
   do i = 0, 32 * 32 - 1
      if (arr(i) .ne. i + (i / 32) * 2) stop 1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-6.f90
index 02e09b3..b2bcf11 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-6.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-6.f90
@@ -3,11 +3,15 @@
 
 ! { dg-do run }
 
+! { dg-additional-options "--param=openacc-kernels=decompose" }
+
 ! { dg-additional-options "-fopt-info-omp-all" }
 ! { dg-additional-options "-foffload=-fopt-info-omp-all" }
 
 ! { dg-additional-options "--param=openacc-privatization=noisy" }
 ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+! Prune a few: uninteresting:
+! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
 
 ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
 ! passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -29,7 +33,8 @@ program main
      arr(i) = i
   end do
 
-  !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] }
+  !$acc kernels copy(arr)
+  ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 }
   !$acc loop gang(num:32) private(pt) ! { dg-line l_loop_i[incr c_loop_i] }
   ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
   ! { dg-note {variable 'pt' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_i$c_loop_i }
@@ -45,8 +50,8 @@ program main
         arr(i * 32 + j) = arr(i * 32 + j) + pt%x + pt%y + pt%z + pt%attr(5);
      end do
   end do
+  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i }
   !$acc end kernels
-  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute }
 
   do i = 0, 32 * 32 - 1
      if (arr(i) .ne. i + (i / 32) * 13) stop 1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-1.f90
index 5811d0c..eba7263 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-1.f90
@@ -2,11 +2,15 @@
 
 ! { dg-do run }
 
+! { dg-additional-options "--param=openacc-kernels=decompose" }
+
 ! { dg-additional-options "-fopt-info-omp-all" }
 ! { dg-additional-options "-foffload=-fopt-info-omp-all" }
 
 ! { dg-additional-options "--param=openacc-privatization=noisy" }
 ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+! Prune a few: uninteresting:
+! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
 
 ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
 ! passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -23,7 +27,8 @@ program main
      arr(i) = i
   end do
 
-  !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] }
+  !$acc kernels copy(arr)
+  ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 }
   !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] }
   ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
   do i = 0, 31
@@ -46,8 +51,8 @@ program main
         end do
      end do
   end do
+  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i }
   !$acc end kernels
-  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute }
 
   do i = 0, 32 - 1
      do j = 0, 32 -1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-2.f90
index 81125a2..7d1c688 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-vector-2.f90
@@ -2,11 +2,15 @@
 
 ! { dg-do run }
 
+! { dg-additional-options "--param=openacc-kernels=decompose" }
+
 ! { dg-additional-options "-fopt-info-omp-all" }
 ! { dg-additional-options "-foffload=-fopt-info-omp-all" }
 
 ! { dg-additional-options "--param=openacc-privatization=noisy" }
 ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+! Prune a few: uninteresting:
+! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
 
 ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
 ! passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -23,7 +27,8 @@ program main
      arr(i) = i
   end do
 
-  !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] }
+  !$acc kernels copy(arr)
+  ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 }
   !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] }
   ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
   do i = 0, 31
@@ -41,8 +46,8 @@ program main
         end do
      end do
   end do
+  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i }
   !$acc end kernels
-  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute }
 
   do i = 0, 32 - 1
      do j = 0, 32 -1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-1.f90
index 824c198..d459a6f 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-1.f90
@@ -2,11 +2,15 @@
 
 ! { dg-do run }
 
+! { dg-additional-options "--param=openacc-kernels=decompose" }
+
 ! { dg-additional-options "-fopt-info-omp-all" }
 ! { dg-additional-options "-foffload=-fopt-info-omp-all" }
 
 ! { dg-additional-options "--param=openacc-privatization=noisy" }
 ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+! Prune a few: uninteresting:
+! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
 
 ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
 ! passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -25,6 +29,8 @@ program main
   end do
 
   !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] }
+  ! { dg-note {variable 'x\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute }
+  ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 }
   !$acc loop gang(num:32) private(x) ! { dg-line l_loop_i[incr c_loop_i] }
   ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
   do i = 0, 31
@@ -36,8 +42,8 @@ program main
         arr(i * 32 + j) = arr(i * 32 + j) + x
      end do
   end do
+  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i }
   !$acc end kernels
-  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute }
 
   do i = 0, 32 * 32 - 1
      if (arr(i) .ne. i + ieor(i / 32, mod(i, 32) * 3)) stop 1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-2.f90
index d25d419..c3a4e2a 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-2.f90
@@ -3,11 +3,15 @@
 
 ! { dg-do run }
 
+! { dg-additional-options "--param=openacc-kernels=decompose" }
+
 ! { dg-additional-options "-fopt-info-omp-all" }
 ! { dg-additional-options "-foffload=-fopt-info-omp-all" }
 
 ! { dg-additional-options "--param=openacc-privatization=noisy" }
 ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+! Prune a few: uninteresting:
+! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
 
 ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
 ! passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -24,7 +28,8 @@ program main
      arr(i) = i
   end do
 
-  !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] }
+  !$acc kernels copy(arr)
+  ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 }
   !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] }
   ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
   do i = 0, 31
@@ -41,8 +46,8 @@ program main
         end do
      end do
   end do
+  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i }
   !$acc end kernels
-  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute }
 
   do i = 0, 32 - 1
      do j = 0, 32 -1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-3.f90
index 7a69145..db6097b 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-3.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-3.f90
@@ -3,11 +3,15 @@
 
 ! { dg-do run }
 
+! { dg-additional-options "--param=openacc-kernels=decompose" }
+
 ! { dg-additional-options "-fopt-info-omp-all" }
 ! { dg-additional-options "-foffload=-fopt-info-omp-all" }
 
 ! { dg-additional-options "--param=openacc-privatization=noisy" }
 ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+! Prune a few: uninteresting:
+! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
 
 ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
 ! passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -24,7 +28,8 @@ program main
      arr(i) = i
   end do
 
-  !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] }
+  !$acc kernels copy(arr)
+  ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 }
   !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] }
   ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
   do i = 0, 31
@@ -54,8 +59,8 @@ program main
         end do
      end do
   end do
+  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i }
   !$acc end kernels
-  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute }
 
   do i = 0, 32 - 1
      do j = 0, 32 -1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-4.f90
index 2c1d566..a0b6c29 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-4.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-4.f90
@@ -3,11 +3,15 @@
 
 ! { dg-do run }
 
+! { dg-additional-options "--param=openacc-kernels=decompose" }
+
 ! { dg-additional-options "-fopt-info-omp-all" }
 ! { dg-additional-options "-foffload=-fopt-info-omp-all" }
 
 ! { dg-additional-options "--param=openacc-privatization=noisy" }
 ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+! Prune a few: uninteresting:
+! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
 
 ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
 ! passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -24,7 +28,8 @@ program main
      arr(i) = i
   end do
 
-  !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] }
+  !$acc kernels copy(arr)
+  ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 }
   !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] }
   ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
   do i = 0, 31
@@ -49,8 +54,8 @@ program main
         end do
      end do
   end do
+  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i }
   !$acc end kernels
-  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute }
 
   do i = 0, 32 - 1
      do j = 0, 32 -1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-5.f90
index 4936e56..22c0b9b 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-5.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-5.f90
@@ -3,11 +3,15 @@
 
 ! { dg-do run }
 
+! { dg-additional-options "--param=openacc-kernels=decompose" }
+
 ! { dg-additional-options "-fopt-info-omp-all" }
 ! { dg-additional-options "-foffload=-fopt-info-omp-all" }
 
 ! { dg-additional-options "--param=openacc-privatization=noisy" }
 ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+! Prune a few: uninteresting:
+! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
 
 ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
 ! passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -27,6 +31,8 @@ program main
   end do
 
   !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] }
+  ! { dg-note {variable 'x\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute }
+  ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 }
   !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] }
   ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
   do i = 0, 31
@@ -53,8 +59,8 @@ program main
         end do
      end do
   end do
+  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i }
   !$acc end kernels
-  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute }
 
   do i = 0, 32 - 1
      do j = 0, 32 -1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-6.f90
index 6b2ec1a..1b4d503 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-6.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-6.f90
@@ -3,11 +3,15 @@
 
 ! { dg-do run }
 
+! { dg-additional-options "--param=openacc-kernels=decompose" }
+
 ! { dg-additional-options "-fopt-info-omp-all" }
 ! { dg-additional-options "-foffload=-fopt-info-omp-all" }
 
 ! { dg-additional-options "--param=openacc-privatization=noisy" }
 ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+! Prune a few: uninteresting:
+! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
 
 ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
 ! passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -29,7 +33,8 @@ program main
      arr(i) = i
   end do
 
-  !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] }
+  !$acc kernels copy(arr)
+  ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 }
   !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] }
   ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
   do i = 0, 31
@@ -53,8 +58,8 @@ program main
         end do
      end do
   end do
+  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i }
   !$acc end kernels
-  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute }
 
   do i = 0, 32 - 1
      do j = 0, 32 -1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-7.f90
index a90be1d..9ee8a5a 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-7.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-worker-7.f90
@@ -3,11 +3,15 @@
 
 ! { dg-do run }
 
+! { dg-additional-options "--param=openacc-kernels=decompose" }
+
 ! { dg-additional-options "-fopt-info-omp-all" }
 ! { dg-additional-options "-foffload=-fopt-info-omp-all" }
 
 ! { dg-additional-options "--param=openacc-privatization=noisy" }
 ! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
+! Prune a few: uninteresting:
+! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
 
 ! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
 ! passed to 'incr' may be unset, and in that case, it will be set to [...]",
@@ -24,7 +28,8 @@ program main
      arr(i) = i
   end do
 
-  !$acc kernels copy(arr) ! { dg-line l_compute[incr c_compute] }
+  !$acc kernels copy(arr)
+  ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 }
   !$acc loop gang(num:32) ! { dg-line l_loop_i[incr c_loop_i] }
   ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i }
   do i = 0, 31
@@ -48,8 +53,8 @@ program main
         end do
      end do
   end do
+  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i }
   !$acc end kernels
-  ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute }
 
   do i = 0, 32 - 1
      do j = 0, 32 -1