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

Message ID 871qzhojdr.fsf@euler.schwinge.homeip.net
State New
Headers
Series Test 'libgomp.oacc-*/kernels-private-vars-*' with '--param=openacc-kernels=decompose' [PR104784] |

Commit Message

Thomas Schwinge March 4, 2022, 2:57 p.m. UTC
  Hi!

On 2022-03-04T14:46:25+0100, I wrote:
> Pushed to master branch commit 8935589b496f755e08cadf26d8ceddf0dd6e0968
> "OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs
> [PR100280, PR104132, PR104133]", see attached.

Pushed to master branch commit e28eb86c18ed765dceb3c56471a848e9f0e120ff
"Test 'libgomp.oacc-*/kernels-private-vars-*' with
'--param=openacc-kernels=decompose' [PR104784]", see attached.


Grüße
 Thomas


-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
  

Patch

From e28eb86c18ed765dceb3c56471a848e9f0e120ff Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 16 Feb 2022 22:24:03 +0100
Subject: [PATCH] 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.
---
 .../kernels-private-vars-local-worker-1.c     | 23 +++++++++++++++----
 .../kernels-private-vars-local-worker-2.c     | 20 ++++++++++++----
 .../kernels-private-vars-local-worker-3.c     | 20 ++++++++++++----
 .../kernels-private-vars-local-worker-4.c     | 20 ++++++++++++----
 .../kernels-private-vars-local-worker-5.c     | 20 ++++++++++++----
 .../kernels-private-vars-loop-gang-1.c        | 11 ++++++---
 .../kernels-private-vars-loop-gang-2.c        | 11 ++++++---
 .../kernels-private-vars-loop-gang-3.c        | 11 ++++++---
 .../kernels-private-vars-loop-gang-4.c        | 10 ++++++--
 .../kernels-private-vars-loop-gang-5.c        | 11 ++++++---
 .../kernels-private-vars-loop-gang-6.c        | 11 ++++++---
 .../kernels-private-vars-loop-vector-1.c      | 20 ++++++++++++----
 .../kernels-private-vars-loop-vector-2.c      | 20 ++++++++++++----
 .../kernels-private-vars-loop-worker-1.c      | 13 +++++++----
 .../kernels-private-vars-loop-worker-2.c      | 20 ++++++++++++----
 .../kernels-private-vars-loop-worker-3.c      | 23 +++++++++++++++----
 .../kernels-private-vars-loop-worker-4.c      | 20 ++++++++++++----
 .../kernels-private-vars-loop-worker-5.c      | 21 +++++++++++++----
 .../kernels-private-vars-loop-worker-6.c      | 20 ++++++++++++----
 .../kernels-private-vars-loop-worker-7.c      | 20 ++++++++++++----
 .../kernels-private-vars-loop-gang-1.f90      |  9 ++++++--
 .../kernels-private-vars-loop-gang-2.f90      |  9 ++++++--
 .../kernels-private-vars-loop-gang-3.f90      |  9 ++++++--
 .../kernels-private-vars-loop-gang-6.f90      |  9 ++++++--
 .../kernels-private-vars-loop-vector-1.f90    |  9 ++++++--
 .../kernels-private-vars-loop-vector-2.f90    |  9 ++++++--
 .../kernels-private-vars-loop-worker-1.f90    |  8 ++++++-
 .../kernels-private-vars-loop-worker-2.f90    |  9 ++++++--
 .../kernels-private-vars-loop-worker-3.f90    |  9 ++++++--
 .../kernels-private-vars-loop-worker-4.f90    |  9 ++++++--
 .../kernels-private-vars-loop-worker-5.f90    |  8 ++++++-
 .../kernels-private-vars-loop-worker-6.f90    |  9 ++++++--
 .../kernels-private-vars-loop-worker-7.f90    |  9 ++++++--
 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 acbeb65273f..43bfaf331ab 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 2558a68eb94..c40c2ab33c5 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 b2a208c163f..bd04dcc7b02 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 46c395620fe..4303ab848ea 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 4b5a15e6ad4..8d0e846ce23 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 4a824941427..470d5e65da4 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 039053f3c86..34513e35f22 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 2b89659a007..ef445c82a93 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 70760705925..0658251ea2a 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 edf0e24af8b..a5118666c63 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 a2df33b767d..119bda3b2dc 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 51c1de53414..5a70bb880e5 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 cb90eaab99d..f5bccabbe6f 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 54e1c93714b..9b2c49ac301 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 80ac99013d6..40baae34562 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 a05ac609123..c8b089c0f59 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 d46bb948626..c1819d27179 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 644c6175863..90955aa8f6e 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 182a12a1b3f..f093cfe630b 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 bdfbb59c5c5..906119caf24 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 09ab3956624..b4f44d69de1 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 bec1069c2a8..f4194c96f90 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 9fde012c19c..da85ac485ad 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 02e09b31cf7..b2bcf11cef0 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 5811d0c41b7..eba7263ae56 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 81125a24d0d..7d1c6885cee 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 824c198799c..d459a6f185e 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 d25d419a11e..c3a4e2a3406 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 7a69145b3b6..db6097b6fd5 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 2c1d5665ce5..a0b6c29d466 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 4936e569b44..22c0b9b1045 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 6b2ec1a0047..1b4d5036596 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 a90be1d45e8..9ee8a5afc2e 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
-- 
2.34.1