[OpenACC privatization] Largely extend diagnostics and corresponding testsuite coverage [PR90115]

	gcc/
	PR middle-end/90115
	* flag-types.h (enum openacc_privatization): New.
	* params.opt (-param=openacc-privatization): New.
	* doc/invoke.texi (openacc-privatization): Document it.
	* omp-general.h (get_openacc_privatization_dump_flags): New
	function.
	* omp-low.c (oacc_privatization_candidate_p): Add diagnostics.
	* omp-offload.c (execute_oacc_device_lower)
	<IFN_UNIQUE_OACC_PRIVATE>: Re-work diagnostics.
	* target.def (goacc.adjust_private_decl): Add 'location_t'
	parameter.
	* doc/tm.texi: Regenerate.
	* config/gcn/gcn-protos.h (gcn_goacc_adjust_private_decl): Adjust.
	* config/gcn/gcn-tree.c (gcn_goacc_adjust_private_decl): Likewise.
	* config/nvptx/nvptx.c (nvptx_goacc_adjust_private_decl):
	Likewise.  Preserve it for...
	(nvptx_goacc_expand_var_decl): ... use here.
	gcc/testsuite/
	PR middle-end/90115
	* c-c++-common/goacc/privatization-1-compute-loop.c: New file.
	* c-c++-common/goacc/privatization-1-compute.c: Likewise.
	* c-c++-common/goacc/privatization-1-routine_gang-loop.c:
	Likewise.
	* c-c++-common/goacc/privatization-1-routine_gang.c: Likewise.
	* gfortran.dg/goacc/privatization-1-compute-loop.f90: Likewise.
	* gfortran.dg/goacc/privatization-1-compute.f90: Likewise.
	* gfortran.dg/goacc/privatization-1-routine_gang-loop.f90:
	Likewise.
	* gfortran.dg/goacc/privatization-1-routine_gang.f90: Likewise.
	* c-c++-common/goacc-gomp/nesting-1.c: Update.
	* c-c++-common/goacc/private-reduction-1.c: Likewise.
	* gfortran.dg/goacc/private-3.f95: Likewise.
	libgomp/
	PR middle-end/90115
	* testsuite/libgomp.oacc-fortran/private-atomic-1-vector.f90: New
	file.
	* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Update.
	* testsuite/libgomp.oacc-c-c++-common/host_data-7.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-1.c:
	Likewise.
	* 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-c-c++-common/loop-g-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-g-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-gwv-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-v-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/private-atomic-1-gang.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/private-atomic-1.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/private-variables.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-variable-1.c:
	Likewise.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise.
	* testsuite/libgomp.oacc-fortran/declare-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/host_data-5.F90: Likewise.
	* testsuite/libgomp.oacc-fortran/if-1.f90: 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.
	* testsuite/libgomp.oacc-fortran/optional-private.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/parallel-dims.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/private-atomic-1-gang.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/private-atomic-1-worker.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/private-variables.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-7.f90: Likewise.

22 files changed, 833 insertions, 35 deletions

diff --git a/gcc/config/gcn/gcn-protos.h b/gcc/config/gcn/gcn-protos.h
index 7ef7ae8..8bd0b43 100644
--- a/gcc/config/gcn/gcn-protos.h
+++ b/gcc/config/gcn/gcn-protos.h
@@ -40,7 +40,7 @@ extern rtx gcn_gen_undef (machine_mode);
 extern bool gcn_global_address_p (rtx);
 extern tree gcn_goacc_adjust_propagation_record (tree record_type, bool sender,
 						 const char *name);
-extern tree gcn_goacc_adjust_private_decl (tree var, int level);
+extern tree gcn_goacc_adjust_private_decl (location_t, tree var, int level);
 extern void gcn_goacc_reduction (gcall *call);
 extern bool gcn_hard_regno_rename_ok (unsigned int from_reg,
 				      unsigned int to_reg);
diff --git a/gcc/config/gcn/gcn-tree.c b/gcc/config/gcn/gcn-tree.c
index 75ea50c..1eb8882 100644
--- a/gcc/config/gcn/gcn-tree.c
+++ b/gcc/config/gcn/gcn-tree.c
@@ -578,7 +578,7 @@ gcn_goacc_adjust_propagation_record (tree record_type, bool sender,
 }
 
 tree
-gcn_goacc_adjust_private_decl (tree var, int level)
+gcn_goacc_adjust_private_decl (location_t, tree var, int level)
 {
   if (level != GOMP_DIM_GANG)
     return var;
diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 60d3f07..6642bdf 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -6680,7 +6680,7 @@ nvptx_truly_noop_truncation (poly_uint64, poly_uint64)
 /* Implement TARGET_GOACC_ADJUST_PRIVATE_DECL.  */
 
 static tree
-nvptx_goacc_adjust_private_decl (tree decl, int level)
+nvptx_goacc_adjust_private_decl (location_t loc, tree decl, int level)
 {
   gcc_checking_assert (!lookup_attribute ("oacc gang-private",
 					  DECL_ATTRIBUTES (decl)));
@@ -6689,14 +6689,12 @@ nvptx_goacc_adjust_private_decl (tree decl, int level)
      declarations.  */
   if (level == GOMP_DIM_GANG)
     {
-      if (dump_file && (dump_flags & TDF_DETAILS))
-	{
-	  fprintf (dump_file, "Setting 'oacc gang-private' attribute for decl:");
-	  print_generic_decl (dump_file, decl, TDF_SLIM);
-	  fputc ('\n', dump_file);
-	}
       tree id = get_identifier ("oacc gang-private");
-      DECL_ATTRIBUTES (decl) = tree_cons (id, NULL, DECL_ATTRIBUTES (decl));
+      /* For later diagnostic purposes, pass LOC as VALUE (wrapped as a
+	 TREE).  */
+      tree loc_tree = build_empty_stmt (loc);
+      DECL_ATTRIBUTES (decl)
+	= tree_cons (id, loc_tree, DECL_ATTRIBUTES (decl));
     }
 
   return decl;
@@ -6708,7 +6706,8 @@ static rtx
 nvptx_goacc_expand_var_decl (tree var)
 {
   /* Place "oacc gang-private" variables in shared memory.  */
-  if (lookup_attribute ("oacc gang-private", DECL_ATTRIBUTES (var)))
+  if (tree attr = lookup_attribute ("oacc gang-private",
+				    DECL_ATTRIBUTES (var)))
     {
       gcc_checking_assert (VAR_P (var));
 
@@ -6728,6 +6727,50 @@ nvptx_goacc_expand_var_decl (tree var)
 	  bool existed = gang_private_shared_hmap.put (var, offset);
 	  gcc_checking_assert (!existed);
 	  gang_private_shared_size += tree_to_uhwi (DECL_SIZE_UNIT (var));
+
+	  location_t loc = EXPR_LOCATION (TREE_VALUE (attr));
+#if 0 /* For some reason, this doesn't work.  */
+	  if (dump_enabled_p ())
+	    {
+	      dump_flags_t l_dump_flags
+		= get_openacc_privatization_dump_flags ();
+
+	      const dump_user_location_t d_u_loc
+		= dump_user_location_t::from_location_t (loc);
+/* PR100695 "Format decoder, quoting in 'dump_printf' etc." */
+#if __GNUC__ >= 10
+# pragma GCC diagnostic push
+# pragma GCC diagnostic ignored "-Wformat"
+#endif
+	      dump_printf_loc (l_dump_flags, d_u_loc,
+			       "variable %<%T%> adjusted for OpenACC"
+			       " privatization level: %qs\n",
+			       var, "gang");
+#if __GNUC__ >= 10
+# pragma GCC diagnostic pop
+#endif
+	    }
+#else /* ..., thus emulate that, good enough for testsuite usage.  */
+	  if (param_openacc_privatization != OPENACC_PRIVATIZATION_QUIET)
+	    inform (loc,
+		    "variable %qD adjusted for OpenACC privatization level:"
+		    " %qs",
+		    var, "gang");
+	  if (dump_file && (dump_flags & TDF_DETAILS))
+	    {
+	      /* 'dumpfile.c:dump_loc' */
+	      fprintf (dump_file, "%s:%d:%d: ", LOCATION_FILE (loc),
+		       LOCATION_LINE (loc), LOCATION_COLUMN (loc));
+	      fprintf (dump_file, "%s: ", "note");
+
+	      fprintf (dump_file,
+		       "variable '");
+	      print_generic_expr (dump_file, var, TDF_SLIM);
+	      fprintf (dump_file,
+		       "' adjusted for OpenACC privatization level: '%s'\n",
+		       "gang");
+	    }
+#endif
 	}
       rtx addr = plus_constant (Pmode, gang_private_shared_sym, offset);
       return gen_rtx_MEM (TYPE_MODE (TREE_TYPE (var)), addr);
diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi
index 9bcbcdc..5cd4e2d 100644
--- a/gcc/doc/invoke.texi
+++ b/gcc/doc/invoke.texi
@@ -14425,6 +14425,14 @@ With @option{--param=openacc-kernels=parloops}, OpenACC `kernels'
 constructs are handled by the @samp{parloops} pass, en bloc.
 This is the current default.
 
+@item openacc-privatization
+Specify mode of OpenACC privatization diagnostics for
+@option{-fopt-info-omp-note} and applicable
+@option{-fdump-tree-*-details}.
+With @option{--param=openacc-privatization=quiet}, don't diagnose.
+This is the current default.
+With @option{--param=openacc-privatization=noisy}, do diagnose.
+
 @end table
 
 The following choices of @var{name} are available on AArch64 targets:
diff --git a/gcc/doc/tm.texi b/gcc/doc/tm.texi
index 78c330c..e3a080e 100644
--- a/gcc/doc/tm.texi
+++ b/gcc/doc/tm.texi
@@ -6236,12 +6236,13 @@ like @code{cond_add@var{m}}.  The default implementation returns a zero
 constant of type @var{type}.
 @end deftypefn
 
-@deftypefn {Target Hook} tree TARGET_GOACC_ADJUST_PRIVATE_DECL (tree @var{var}, int @var{level})
+@deftypefn {Target Hook} tree TARGET_GOACC_ADJUST_PRIVATE_DECL (location_t @var{loc}, tree @var{var}, int @var{level})
 This hook, if defined, is used by accelerator target back-ends to adjust
 OpenACC variable declarations that should be made private to the given
 parallelism level (i.e. @code{GOMP_DIM_GANG}, @code{GOMP_DIM_WORKER} or
 @code{GOMP_DIM_VECTOR}).  A typical use for this hook is to force variable
 declarations at the @code{gang} level to reside in GPU shared memory.
+@var{loc} may be used for diagnostic purposes.
 
 You may also use the @code{TARGET_GOACC_EXPAND_VAR_DECL} hook if the
 adjusted variable declaration needs to be expanded to RTL in a non-standard
diff --git a/gcc/flag-types.h b/gcc/flag-types.h
index d60bb30..375448e 100644
--- a/gcc/flag-types.h
+++ b/gcc/flag-types.h
@@ -442,6 +442,13 @@ enum openacc_kernels
   OPENACC_KERNELS_PARLOOPS
 };
 
+/* Modes of OpenACC privatization diagnostics.  */
+enum openacc_privatization
+{
+  OPENACC_PRIVATIZATION_QUIET,
+  OPENACC_PRIVATIZATION_NOISY
+};
+
 #endif
 
 #endif /* ! GCC_FLAG_TYPES_H */
diff --git a/gcc/omp-general.h b/gcc/omp-general.h
index aa04895..5c3e0f0 100644
--- a/gcc/omp-general.h
+++ b/gcc/omp-general.h
@@ -132,4 +132,17 @@ enum omp_requires {
 
 extern GTY(()) enum omp_requires omp_requires_mask;
 
+static inline dump_flags_t
+get_openacc_privatization_dump_flags ()
+{
+  dump_flags_t l_dump_flags = MSG_NOTE;
+
+  /* For '--param=openacc-privatization=quiet', diagnostics only go to dump
+     files.  */
+  if (param_openacc_privatization == OPENACC_PRIVATIZATION_QUIET)
+    l_dump_flags |= MSG_PRIORITY_INTERNALS;
+
+  return l_dump_flags;
+}
+
 #endif /* GCC_OMP_GENERAL_H */
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 577676b..0d63e82 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -10160,16 +10160,81 @@ lower_omp_for_lastprivate (struct omp_for_data *fd, gimple_seq *body_p,
    sometimes, using shared memory directly would be faster than
    broadcasting.  */
 
+static void
+oacc_privatization_begin_diagnose_var (const dump_flags_t l_dump_flags,
+				       const location_t loc, const tree c,
+				       const tree decl)
+{
+  const dump_user_location_t d_u_loc
+    = dump_user_location_t::from_location_t (loc);
+/* PR100695 "Format decoder, quoting in 'dump_printf' etc." */
+#if __GNUC__ >= 10
+# pragma GCC diagnostic push
+# pragma GCC diagnostic ignored "-Wformat"
+#endif
+  dump_printf_loc (l_dump_flags, d_u_loc,
+		   "variable %<%T%> ", decl);
+#if __GNUC__ >= 10
+# pragma GCC diagnostic pop
+#endif
+  if (c)
+    dump_printf (l_dump_flags,
+		 "in %qs clause ",
+		 omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+  else
+    dump_printf (l_dump_flags,
+		 "declared in block ");
+}
+
 static bool
-oacc_privatization_candidate_p (const tree decl)
+oacc_privatization_candidate_p (const location_t loc, const tree c,
+				const tree decl)
 {
+  dump_flags_t l_dump_flags = get_openacc_privatization_dump_flags ();
+
   bool res = true;
 
   if (res && !VAR_P (decl))
-    res = false;
+    {
+      res = false;
+
+      if (dump_enabled_p ())
+	{
+	  oacc_privatization_begin_diagnose_var (l_dump_flags, loc, c, decl);
+	  dump_printf (l_dump_flags,
+		       "potentially has improper OpenACC privatization level: %qs\n",
+		       get_tree_code_name (TREE_CODE (decl)));
+	}
+    }
 
   if (res && !TREE_ADDRESSABLE (decl))
-    res = false;
+    {
+      res = false;
+
+      if (dump_enabled_p ())
+	{
+	  oacc_privatization_begin_diagnose_var (l_dump_flags, loc, c, decl);
+	  dump_printf (l_dump_flags,
+		       "isn%'t candidate for adjusting OpenACC privatization level: %s\n",
+		       "not addressable");
+	}
+    }
+
+  if (res)
+    {
+      if (dump_enabled_p ())
+	{
+	  oacc_privatization_begin_diagnose_var (l_dump_flags, loc, c, decl);
+	  dump_printf (l_dump_flags,
+		       "is candidate for adjusting OpenACC privatization level\n");
+	}
+    }
+
+  if (dump_file && (dump_flags & TDF_DETAILS))
+    {
+      print_generic_decl (dump_file, decl, dump_flags);
+      fprintf (dump_file, "\n");
+    }
 
   return res;
 }
@@ -10185,7 +10250,7 @@ oacc_privatization_scan_clause_chain (omp_context *ctx, tree clauses)
       {
 	tree decl = OMP_CLAUSE_DECL (c);
 
-	if (!oacc_privatization_candidate_p (decl))
+	if (!oacc_privatization_candidate_p (OMP_CLAUSE_LOCATION (c), c, decl))
 	  continue;
 
 	gcc_checking_assert (!ctx->oacc_privatization_candidates.contains (decl));
@@ -10201,7 +10266,7 @@ oacc_privatization_scan_decl_chain (omp_context *ctx, tree decls)
 {
   for (tree decl = decls; decl; decl = DECL_CHAIN (decl))
     {
-      if (!oacc_privatization_candidate_p (decl))
+      if (!oacc_privatization_candidate_p (gimple_location (ctx->stmt), NULL, decl))
 	continue;
 
       gcc_checking_assert (!ctx->oacc_privatization_candidates.contains (decl));
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 8bfb8b3..e907827 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -2137,6 +2137,15 @@ execute_oacc_device_lower ()
 
 		case IFN_UNIQUE_OACC_PRIVATE:
 		  {
+		    dump_flags_t l_dump_flags
+		      = get_openacc_privatization_dump_flags ();
+
+		    location_t loc = gimple_location (stmt);
+		    if (LOCATION_LOCUS (loc) == UNKNOWN_LOCATION)
+		      loc = DECL_SOURCE_LOCATION (current_function_decl);
+		    const dump_user_location_t d_u_loc
+		      = dump_user_location_t::from_location_t (loc);
+
 		    HOST_WIDE_INT level
 		      = TREE_INT_CST_LOW (gimple_call_arg (call, 2));
 		    gcc_checking_assert (level == -1
@@ -2146,31 +2155,65 @@ execute_oacc_device_lower ()
 			 i < gimple_call_num_args (call);
 			 i++)
 		      {
+			static char const *const axes[] =
+			/* Must be kept in sync with GOMP_DIM enumeration.  */
+			  { "gang", "worker", "vector" };
+
 			tree arg = gimple_call_arg (call, i);
 			gcc_checking_assert (TREE_CODE (arg) == ADDR_EXPR);
 			tree decl = TREE_OPERAND (arg, 0);
-			if (dump_file && (dump_flags & TDF_DETAILS))
+			if (dump_enabled_p ())
+/* PR100695 "Format decoder, quoting in 'dump_printf' etc." */
+#if __GNUC__ >= 10
+# pragma GCC diagnostic push
+# pragma GCC diagnostic ignored "-Wformat"
+#endif
+			  dump_printf_loc (l_dump_flags, d_u_loc,
+					   "variable %<%T%> ought to be"
+					   " adjusted for OpenACC"
+					   " privatization level: %qs\n",
+					   decl,
+					   (level == -1
+					    ? "UNKNOWN" : axes[level]));
+#if __GNUC__ >= 10
+# pragma GCC diagnostic pop
+#endif
+			bool adjusted;
+			if (level == -1)
+			  adjusted = false;
+			else if (!targetm.goacc.adjust_private_decl)
+			  adjusted = false;
+			else if (level == GOMP_DIM_VECTOR)
 			  {
-			    static char const *const axes[] =
-			      /* Must be kept in sync with GOMP_DIM
-				 enumeration.  */
-			      { "gang", "worker", "vector" };
-			    fprintf (dump_file, "Decl UID %u has %s "
-				     "partitioning:", DECL_UID (decl),
-				     (level == -1 ? "UNKNOWN" : axes[level]));
-			    print_generic_decl (dump_file, decl, TDF_SLIM);
-			    fputc ('\n', dump_file);
+			    /* That's the default behavior.  */
+			    adjusted = true;
 			  }
-			if (level != -1
-			    && targetm.goacc.adjust_private_decl)
+			else
 			  {
 			    tree oldtype = TREE_TYPE (decl);
 			    tree newdecl
-			      = targetm.goacc.adjust_private_decl (decl, level);
-			    if (TREE_TYPE (newdecl) != oldtype
-				|| newdecl != decl)
+			      = targetm.goacc.adjust_private_decl (loc, decl,
+								   level);
+			    adjusted = (TREE_TYPE (newdecl) != oldtype
+					|| newdecl != decl);
+			    if (adjusted)
 			      adjusted_vars.put (decl, newdecl);
 			  }
+			if (adjusted
+			    && dump_enabled_p ())
+/* PR100695 "Format decoder, quoting in 'dump_printf' etc." */
+#if __GNUC__ >= 10
+# pragma GCC diagnostic push
+# pragma GCC diagnostic ignored "-Wformat"
+#endif
+			  dump_printf_loc (l_dump_flags, d_u_loc,
+					   "variable %<%T%> adjusted for"
+					   " OpenACC privatization level:"
+					   " %qs\n",
+					   decl, axes[level]);
+#if __GNUC__ >= 10
+# pragma GCC diagnostic pop
+#endif
 		      }
 		    remove = true;
 		  }
diff --git a/gcc/params.opt b/gcc/params.opt
index 82600b9..0d0dcd2 100644
--- a/gcc/params.opt
+++ b/gcc/params.opt
@@ -795,6 +795,19 @@ Enum(openacc_kernels) String(decompose) Value(OPENACC_KERNELS_DECOMPOSE)
 EnumValue
 Enum(openacc_kernels) String(parloops) Value(OPENACC_KERNELS_PARLOOPS)
 
+-param=openacc-privatization=
+Common Joined Enum(openacc_privatization) Var(param_openacc_privatization) Init(OPENACC_PRIVATIZATION_QUIET) Param
+--param=openacc-privatization=[quiet|noisy]	Specify mode of OpenACC privatization diagnostics.
+
+Enum
+Name(openacc_privatization) Type(enum openacc_privatization)
+
+EnumValue
+Enum(openacc_privatization) String(quiet) Value(OPENACC_PRIVATIZATION_QUIET)
+
+EnumValue
+Enum(openacc_privatization) String(noisy) Value(OPENACC_PRIVATIZATION_NOISY)
+
 -param=parloops-chunk-size=
 Common Joined UInteger Var(param_parloops_chunk_size) Param Optimization
 Chunk size of omp schedule for loops parallelized by parloops.
diff --git a/gcc/target.def b/gcc/target.def
index 660b69f..1dffedc 100644
--- a/gcc/target.def
+++ b/gcc/target.def
@@ -1733,11 +1733,12 @@ OpenACC variable declarations that should be made private to the given\n\
 parallelism level (i.e. @code{GOMP_DIM_GANG}, @code{GOMP_DIM_WORKER} or\n\
 @code{GOMP_DIM_VECTOR}).  A typical use for this hook is to force variable\n\
 declarations at the @code{gang} level to reside in GPU shared memory.\n\
+@var{loc} may be used for diagnostic purposes.\n\
 \n\
 You may also use the @code{TARGET_GOACC_EXPAND_VAR_DECL} hook if the\n\
 adjusted variable declaration needs to be expanded to RTL in a non-standard\n\
 way.",
-tree, (tree var, int level),
+tree, (location_t loc, tree var, int level),
 NULL)
 
 DEFHOOK
diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
index aaf0e7a..b0b7837 100644
--- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
+++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
@@ -1,7 +1,14 @@
+/* { dg-additional-options "-fopt-info-omp-note" } */
+/* { dg-additional-options "--param=openacc-privatization=noisy" } for
+   testing/documenting aspects of that functionality.  */
+
+
 void
 f_acc_data (void)
 {
 #pragma acc data
+  /* { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } */
+  /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } .-2 } */
   {
     int i;
 #pragma omp atomic write
@@ -13,6 +20,8 @@ void
 f_acc_kernels (void)
 {
 #pragma acc kernels
+  /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } .-1 }
+     { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } .-2 } */
   {
     int i;
 #pragma omp atomic write
@@ -27,6 +36,9 @@ f_acc_loop (void)
   int i;
 
 #pragma acc loop
+  /* { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } */
+  /* { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } .-2 }
+     { dg-bogus {note: variable 'i' ought to be adjusted for OpenACC privatization level: 'UNKNOWN'} "TODO" { xfail *-*-* } .-3 } */
   for (i = 0; i < 2; ++i)
     {
 #pragma omp atomic write
@@ -38,6 +50,8 @@ void
 f_acc_parallel (void)
 {
 #pragma acc parallel
+  /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } .-1 }
+     { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } .-2 } */
   {
     int i;
 #pragma omp atomic write
diff --git a/gcc/testsuite/c-c++-common/goacc/private-reduction-1.c b/gcc/testsuite/c-c++-common/goacc/private-reduction-1.c
index d4e3995..38f6b7a 100644
--- a/gcc/testsuite/c-c++-common/goacc/private-reduction-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/private-reduction-1.c
@@ -1,3 +1,7 @@
+/* { dg-additional-options "-fopt-info-note-omp" }
+   { dg-additional-options "--param=openacc-privatization=noisy" } for
+   testing/documenting aspects of that functionality.  */
+
 int
 reduction ()
 {
@@ -5,6 +9,8 @@ reduction ()
 
   #pragma acc parallel
   #pragma acc loop private (r) reduction (+:r)
+  /* { dg-note {variable 'r' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} r { target *-*-* } .-1 } */
+  /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} i { target *-*-* } .-2 } */
   for (i = 0; i < 100; i++)
     r += 10;
 
diff --git a/gcc/testsuite/c-c++-common/goacc/privatization-1-compute-loop.c b/gcc/testsuite/c-c++-common/goacc/privatization-1-compute-loop.c
new file mode 100644
index 0000000..4bfb527
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/privatization-1-compute-loop.c
@@ -0,0 +1,95 @@
+/* OpenACC privatization: 'loop' construct inside compute construct */
+
+/* { dg-additional-options "-fopt-info-omp-note" } */
+/* { dg-additional-options "--param=openacc-privatization=noisy" } for
+   testing/documenting aspects of that functionality.  */
+
+/* See also '../../gfortran.dg/goacc/privatization-1-compute-loop.f90'.  */
+
+/* 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 [...]",
+   so to maintain compatibility with earlier Tcl releases, we manually
+   initialize counter variables:
+   { dg-line l_dummy[variable c_loop 0] }
+   { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
+   "WARNING: dg-line var l_dummy defined, but not used".  */
+
+extern int e;
+static int s;
+int g;
+
+void
+f (int i, int j, int a)
+{
+  extern int ex;
+  static int st;
+  int x, y;
+#pragma acc parallel
+#pragma acc loop collapse(2) private(a) private (e, s, g) private(ex, st, x, y) /* { dg-line l_loop[incr c_loop] } */
+  for (i = 0; i < 20; ++i)
+    for (j = 0; j < 25; ++j)
+      {
+	__label__ ll;
+	/* Nested scopes fun.  */
+	{
+	  struct s_ss { int i; } ss;
+	  {
+	    extern int func (int *, int *, int *);
+	    /* Don't know how to effect a 'CONST_DECL' here.  (See Fortran example.)  */
+	    /* Don't know how to effect a 'RESULT_DECL' here; only saw this for OpenMP 'lastprivate'.  */
+
+	    a = func (&i, &j, &a);
+	  }
+	  ss.i = a;
+	  {
+	    extern int func2 (int *, int *, int *, int *, int *, int *, int *);
+	    extern int ext;
+	    static int sta;
+	    a = func2 (&e, &s, &g, &ex, &st, &ext, &sta);
+	  }
+	}
+	x = a;
+#pragma acc atomic write
+	y = a;
+	{
+	  int xx, yy;
+	  xx = a;
+#pragma acc atomic write
+	  yy = a;
+	}
+
+      ll:
+	;
+      }
+  /* { dg-note {variable 'y' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'y' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'st' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'st' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'ex' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'ex' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'g' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'g' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 's' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 's' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'e' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'e' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'a' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'j\.1' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'j' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'i' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'll' declared in block potentially has improper OpenACC privatization level: 'label_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'struct struct s_ss' declared in block potentially has improper OpenACC privatization level: 'type_decl'} "TODO" { target c } l_loop$c_loop }
+     { dg-note {variable 's_ss' declared in block potentially has improper OpenACC privatization level: 'type_decl'} "TODO" { target c++ } l_loop$c_loop } */
+  /* { dg-note {variable 'ss' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'func' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'func2' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'ext' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'ext' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'sta' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'sta' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'xx' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'yy' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'yy' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/privatization-1-compute.c b/gcc/testsuite/c-c++-common/goacc/privatization-1-compute.c
new file mode 100644
index 0000000..4de45e5
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/privatization-1-compute.c
@@ -0,0 +1,90 @@
+/* OpenACC privatization: compute construct */
+
+/* { dg-additional-options "-fopt-info-omp-note" } */
+/* { dg-additional-options "--param=openacc-privatization=noisy" } for
+   testing/documenting aspects of that functionality.  */
+
+/* See also '../../gfortran.dg/goacc/privatization-1-compute.f90'.  */
+
+/* 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 [...]",
+   so to maintain compatibility with earlier Tcl releases, we manually
+   initialize counter variables:
+   { dg-line l_dummy[variable c_compute 0] }
+   { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
+   "WARNING: dg-line var l_dummy defined, but not used".  */
+
+extern int e;
+static int s;
+int g;
+
+void
+f (int i, int j, int a)
+{
+  extern int ex;
+  static int st;
+  int x, y;
+#pragma acc parallel private(i, j, a) private (e, s, g) private(ex, st, x, y) /* { dg-line l_compute[incr c_compute] } */
+      {
+	__label__ ll;
+	/* Nested scopes fun.  */
+	{
+	  struct s_ss { int i; } ss;
+	  {
+	    extern int func (int *, int *, int *);
+	    /* Don't know how to effect a 'CONST_DECL' here.  (See Fortran example.)  */
+	    /* Don't know how to effect a 'RESULT_DECL' here; only saw this for OpenMP 'lastprivate'.  */
+
+	    a = func (&i, &j, &a);
+	  }
+	  ss.i = a;
+	  {
+	    extern int func2 (int *, int *, int *, int *, int *, int *, int *);
+	    extern int ext;
+	    static int sta;
+	    a = func2 (&e, &s, &g, &ex, &st, &ext, &sta);
+	  }
+	}
+	x = a;
+#pragma acc atomic write
+	y = a;
+	{
+	  int xx, yy;
+	  xx = a;
+#pragma acc atomic write
+	  yy = a;
+	}
+
+      ll:
+	;
+      }
+  /* { dg-note {variable 'y' in 'private' clause is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_compute$c_compute }
+     { dg-note {variable 'y' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "TODO" { xfail *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'st' in 'private' clause is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_compute$c_compute }
+     { dg-note {variable 'st' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'ex' in 'private' clause is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_compute$c_compute }
+     { dg-note {variable 'ex' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'g' in 'private' clause is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_compute$c_compute }
+     { dg-note {variable 'g' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 's' in 'private' clause is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_compute$c_compute }
+     { dg-note {variable 's' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'e' in 'private' clause is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_compute$c_compute }
+     { dg-note {variable 'e' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'a' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { xfail *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'j' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { xfail *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'i' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { xfail *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'll' declared in block potentially has improper OpenACC privatization level: 'label_decl'} "TODO" { target *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'struct struct s_ss' declared in block potentially has improper OpenACC privatization level: 'type_decl'} "TODO" { target c } l_compute$c_compute }
+     { dg-note {variable 's_ss' declared in block potentially has improper OpenACC privatization level: 'type_decl'} "TODO" { target c++ } l_compute$c_compute } */
+  /* { dg-note {variable 'ss' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'func' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { target *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'func2' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { target *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'ext' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
+     { dg-note {variable 'ext' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'sta' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
+     { dg-note {variable 'sta' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'xx' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
+  /* { dg-note {variable 'yy' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
+     { dg-note {variable 'yy' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute } */
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/privatization-1-routine_gang-loop.c b/gcc/testsuite/c-c++-common/goacc/privatization-1-routine_gang-loop.c
new file mode 100644
index 0000000..fcc233b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/privatization-1-routine_gang-loop.c
@@ -0,0 +1,95 @@
+/* OpenACC privatization: 'loop' construct inside 'routine' */
+
+/* { dg-additional-options "-fopt-info-omp-note" } */
+/* { dg-additional-options "--param=openacc-privatization=noisy" } for
+   testing/documenting aspects of that functionality.  */
+
+/* See also '../../gfortran.dg/goacc/privatization-1-routine_gang-loop.f90'.  */
+
+/* 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 [...]",
+   so to maintain compatibility with earlier Tcl releases, we manually
+   initialize counter variables:
+   { dg-line l_dummy[variable c_loop 0] }
+   { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
+   "WARNING: dg-line var l_dummy defined, but not used".  */
+
+extern int e;
+static int s;
+int g;
+
+#pragma acc routine gang
+void
+f (int i, int j, int a)
+{
+  extern int ex;
+  static int st;
+  int x, y;
+#pragma acc loop collapse(2) private(a) private (e, s, g) private(ex, st, x, y) /* { dg-line l_loop[incr c_loop] } */
+  for (i = 0; i < 20; ++i)
+    for (j = 0; j < 25; ++j)
+      {
+	__label__ ll;
+	/* Nested scopes fun.  */
+	{
+	  struct s_ss { int i; } ss;
+	  {
+	    extern int func (int *, int *, int *);
+	    /* Don't know how to effect a 'CONST_DECL' here.  (See Fortran example.)  */
+	    /* Don't know how to effect a 'RESULT_DECL' here; only saw this for OpenMP 'lastprivate'.  */
+
+	    a = func (&i, &j, &a);
+	  }
+	  ss.i = a;
+	  {
+	    extern int func2 (int *, int *, int *, int *, int *, int *, int *);
+	    extern int ext;
+	    static int sta;
+	    a = func2 (&e, &s, &g, &ex, &st, &ext, &sta);
+	  }
+	}
+	x = a;
+#pragma acc atomic write
+	y = a;
+	{
+	  int xx, yy;
+	  xx = a;
+#pragma acc atomic write
+	  yy = a;
+	}
+
+      ll:
+	;
+      }
+  /* { dg-note {variable 'y' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'y' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'st' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'st' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'ex' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'ex' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'g' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'g' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 's' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 's' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'e' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'e' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'a' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'j\.1' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'j' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'i' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'll' declared in block potentially has improper OpenACC privatization level: 'label_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'struct struct s_ss' declared in block potentially has improper OpenACC privatization level: 'type_decl'} "TODO" { target c } l_loop$c_loop }
+     { dg-note {variable 's_ss' declared in block potentially has improper OpenACC privatization level: 'type_decl'} "TODO" { target c++ } l_loop$c_loop } */
+  /* { dg-note {variable 'ss' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'func' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'func2' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'ext' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'ext' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'sta' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'sta' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'xx' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
+  /* { dg-note {variable 'yy' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+     { dg-note {variable 'yy' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop } */
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/privatization-1-routine_gang.c b/gcc/testsuite/c-c++-common/goacc/privatization-1-routine_gang.c
new file mode 100644
index 0000000..cd6708f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/privatization-1-routine_gang.c
@@ -0,0 +1,93 @@
+/* OpenACC privatization: 'routine' */
+
+/* { dg-additional-options "-fopt-info-omp-note" } */
+/* { dg-additional-options "--param=openacc-privatization=noisy" } for
+   testing/documenting aspects of that functionality.  */
+
+/* See also '../../gfortran.dg/goacc/privatization-1-routine_gang.f90'.  */
+
+/* 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 [...]",
+   so to maintain compatibility with earlier Tcl releases, we manually
+   initialize counter variables:
+   { dg-line l_dummy[variable c_routine 0] }
+   { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
+   "WARNING: dg-line var l_dummy defined, but not used".  */
+
+extern int e;
+static int s;
+int g;
+#pragma acc declare device_resident(e, s, g)
+
+#pragma acc routine gang /* { dg-line l_routine[incr c_routine] } */
+void
+f (int i, int j, int a)
+{
+  extern int ex;
+  static int st;
+#pragma acc declare device_resident(ex /* , st */)
+  int x, y;
+      {
+	__label__ ll;
+	/* Nested scopes fun.  */
+	{
+	  struct s_ss { int i; } ss;
+	  {
+	    extern int func (int *, int *, int *);
+	    /* Don't know how to effect a 'CONST_DECL' here.  (See Fortran example.)  */
+	    /* Don't know how to effect a 'RESULT_DECL' here; only saw this for OpenMP 'lastprivate'.  */
+
+	    a = func (&i, &j, &a);
+	  }
+	  ss.i = a;
+	  {
+	    extern int func2 (int *, int *, int *, int *, int *, int *, int *);
+	    extern int ext;
+	    static int sta;
+#pragma acc declare device_resident(ext /* , sta */)
+	    a = func2 (&e, &s, &g, &ex, &st, &ext, &sta);
+	  }
+	}
+	x = a;
+#pragma acc atomic write
+	y = a;
+	{
+	  int xx, yy;
+	  xx = a;
+#pragma acc atomic write
+	  yy = a;
+	}
+
+      ll:
+	;
+      }
+}
+  /* { dg-note {variable 'y' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_routine$c_routine }
+     { dg-note {variable 'y' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'st' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_routine$c_routine }
+     { dg-note {variable 'st' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'ex' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_routine$c_routine }
+     { dg-note {variable 'ex' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'g' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_routine$c_routine }
+     { dg-note {variable 'g' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 's' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_routine$c_routine }
+     { dg-note {variable 's' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'e' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_routine$c_routine }
+     { dg-note {variable 'e' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'a' declared in block potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'j' declared in block potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'i' declared in block potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'll' declared in block potentially has improper OpenACC privatization level: 'label_decl'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'struct struct s_ss' declared in block potentially has improper OpenACC privatization level: 'type_decl'} "TODO" { target c xfail *-*-* } l_routine$c_routine }
+     { dg-note {variable 's_ss' declared in block potentially has improper OpenACC privatization level: 'type_decl'} "TODO" { target c++ xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'ss' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'func' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'func2' declared in block potentially has improper OpenACC privatization level: 'function_decl'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'ext' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_routine$c_routine }
+     { dg-note {variable 'ext' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'sta' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_routine$c_routine }
+     { dg-note {variable 'sta' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'xx' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "TODO" { xfail *-*-* } l_routine$c_routine } */
+  /* { dg-note {variable 'yy' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { xfail *-*-* } l_routine$c_routine }
+     { dg-note {variable 'yy' ought to be adjusted for OpenACC privatization level: 'gang'} "TODO" { xfail *-*-* } l_routine$c_routine } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-3.f95 b/gcc/testsuite/gfortran.dg/goacc/private-3.f95
index a7c6d81..1bfb4f1 100644
--- a/gcc/testsuite/gfortran.dg/goacc/private-3.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/private-3.f95
@@ -1,7 +1,9 @@
-! { dg-do compile }
-
 ! test for private variables in a reduction clause
 
+! { dg-additional-options "-fopt-info-note-omp" }
+! { dg-additional-options "--param=openacc-privatization=noisy" } for
+! testing/documenting aspects of that functionality.
+
 program test
   implicit none
   integer, parameter :: n = 100
@@ -16,6 +18,7 @@ program test
   !$acc parallel private (k)
   k = 0
   !$acc loop reduction (+:k)
+  ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
   do i = 1, n
      k = k + 1
   end do
diff --git a/gcc/testsuite/gfortran.dg/goacc/privatization-1-compute-loop.f90 b/gcc/testsuite/gfortran.dg/goacc/privatization-1-compute-loop.f90
new file mode 100644
index 0000000..bcd7159
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/privatization-1-compute-loop.f90
@@ -0,0 +1,57 @@
+! OpenACC privatization: 'loop' construct
+
+! { dg-additional-options "-fopt-info-omp-note" }
+! { dg-additional-options "--param=openacc-privatization=noisy" } for
+! testing/documenting aspects of that functionality.
+
+! See also '../../c-c++-common/goacc/privatization-1-compute-loop.c'.
+!TODO More cases should be added here.
+
+! 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 [...]",
+! so to maintain compatibility with earlier Tcl releases, we manually
+! initialize counter variables:
+! { dg-line l_dummy[variable c_loop 0] }
+! { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
+! "WARNING: dg-line var l_dummy defined, but not used".
+
+module m
+contains
+  subroutine f (i, j, a)
+    implicit none
+    integer :: i, j, a
+    integer :: x, y
+    integer, parameter :: c = 3
+    integer, external :: g
+
+    !$acc parallel
+    !$acc loop collapse(2) private(a) private(x, y) ! { dg-line l_loop[incr c_loop] }
+    do i = 1, 20
+       do j = 1, 25
+          ! Can't have nested scopes fun.  (Fortran 'block' construct supported only starting with OpenACC 3.1.)
+
+          ! Don't know how to effect a 'LABEL_DECL' here.
+          ! Don't know how to effect a 'TYPE_DECL' here.
+          ! Don't know how to effect a 'FUNCTION_DECL' here.
+          ! Don't know how to effect a 'RESULT_DECL' here.
+          ! Don't know how to effect a 'VAR_DECL' here.
+          ! (See C/C++ example.)
+
+          a = g (i, j, a, c)
+          x = a
+          !$acc atomic write
+          y = a
+       end do
+    end do
+    ! { dg-note {variable 'count\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'i' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'j' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'a' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'y' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'y' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop }
+    !$acc end parallel
+  end subroutine f
+end module m
diff --git a/gcc/testsuite/gfortran.dg/goacc/privatization-1-compute.f90 b/gcc/testsuite/gfortran.dg/goacc/privatization-1-compute.f90
new file mode 100644
index 0000000..ed7e9ec
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/privatization-1-compute.f90
@@ -0,0 +1,48 @@
+! OpenACC privatization: compute construct
+
+! { dg-additional-options "-fopt-info-omp-note" }
+! { dg-additional-options "--param=openacc-privatization=noisy" } for
+! testing/documenting aspects of that functionality.
+
+! See also '../../c-c++-common/goacc/privatization-1-compute.c'.
+!TODO More cases should be added here.
+
+! 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 [...]",
+! so to maintain compatibility with earlier Tcl releases, we manually
+! initialize counter variables:
+! { dg-line l_dummy[variable c_compute 0] }
+! { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
+! "WARNING: dg-line var l_dummy defined, but not used".
+
+module m
+contains
+  subroutine f (i, j, a)
+    implicit none
+    integer :: i, j, a
+    integer :: x, y
+    integer, parameter :: c = 3
+    integer, external :: g
+
+    !$acc parallel private(i, j, a) private(x, y) ! { dg-line l_compute[incr c_compute] }
+          ! Can't have nested scopes fun.  (Fortran 'block' construct supported only starting with OpenACC 3.1.)
+
+          ! Don't know how to effect a 'LABEL_DECL' here.
+          ! Don't know how to effect a 'TYPE_DECL' here.
+          ! Don't know how to effect a 'FUNCTION_DECL' here.
+          ! Don't know how to effect a 'RESULT_DECL' here.
+          ! Don't know how to effect a 'VAR_DECL' here.
+          ! (See C/C++ example.)
+
+          a = g (i, j, a, c)
+          x = a
+          !$acc atomic write ! ... to force 'TREE_ADDRESSABLE'.
+          y = a
+    !$acc end parallel
+    ! { dg-note {variable 'i' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { xfail *-*-* } l_compute$c_compute }
+    ! { dg-note {variable 'j' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { xfail *-*-* } l_compute$c_compute }
+    ! { dg-note {variable 'a' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { xfail *-*-* } l_compute$c_compute }
+    ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute }
+    ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute }
+  end subroutine f
+end module m
diff --git a/gcc/testsuite/gfortran.dg/goacc/privatization-1-routine_gang-loop.f90 b/gcc/testsuite/gfortran.dg/goacc/privatization-1-routine_gang-loop.f90
new file mode 100644
index 0000000..db6d822
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/privatization-1-routine_gang-loop.f90
@@ -0,0 +1,56 @@
+! OpenACC privatization: 'loop' construct inside 'routine'
+
+! { dg-additional-options "-fopt-info-omp-note" }
+! { dg-additional-options "--param=openacc-privatization=noisy" } for
+! testing/documenting aspects of that functionality.
+
+! See also '../../c-c++-common/goacc/privatization-1-routine_gang-loop.c'.
+!TODO More cases should be added here.
+
+! 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 [...]",
+! so to maintain compatibility with earlier Tcl releases, we manually
+! initialize counter variables:
+! { dg-line l_dummy[variable c_loop 0] }
+! { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
+! "WARNING: dg-line var l_dummy defined, but not used".
+
+module m
+contains
+  subroutine f (i, j, a)
+    implicit none
+    integer :: i, j, a
+    !$acc routine (f) gang
+    integer :: x, y
+    integer, parameter :: c = 3
+    integer, external :: g
+
+    !$acc loop collapse(2) private(a) private(x, y) ! { dg-line l_loop[incr c_loop] }
+    do i = 1, 20
+       do j = 1, 25
+          ! Can't have nested scopes fun.  (Fortran 'block' construct supported only starting with OpenACC 3.1.)
+
+          ! Don't know how to effect a 'LABEL_DECL' here.
+          ! Don't know how to effect a 'TYPE_DECL' here.
+          ! Don't know how to effect a 'FUNCTION_DECL' here.
+          ! Don't know how to effect a 'RESULT_DECL' here.
+          ! Don't know how to effect a 'VAR_DECL' here.
+          ! (See C/C++ example.)
+
+          a = g (i, j, a, c)
+          x = a
+          !$acc atomic write
+          y = a
+       end do
+    end do
+    ! { dg-note {variable 'count\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'i' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'j' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'a' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'y' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop }
+    ! { dg-note {variable 'y' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop }
+  end subroutine f
+end module m
diff --git a/gcc/testsuite/gfortran.dg/goacc/privatization-1-routine_gang.f90 b/gcc/testsuite/gfortran.dg/goacc/privatization-1-routine_gang.f90
new file mode 100644
index 0000000..725bd5e
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/privatization-1-routine_gang.f90
@@ -0,0 +1,47 @@
+! OpenACC privatization: 'routine'
+
+! { dg-additional-options "-fopt-info-omp-note" }
+! { dg-additional-options "--param=openacc-privatization=noisy" } for
+! testing/documenting aspects of that functionality.
+
+! See also '../../c-c++-common/goacc/privatization-1-routine_gang.c'.
+!TODO More cases should be added here.
+
+! 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 [...]",
+! so to maintain compatibility with earlier Tcl releases, we manually
+! initialize counter variables:
+! { dg-line l_dummy[variable c_routine 0] }
+! { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
+! "WARNING: dg-line var l_dummy defined, but not used".
+
+module m
+contains
+  subroutine f (i, j, a)
+    implicit none
+    integer :: i, j, a
+    !$acc routine (f) gang ! { dg-line l_routine[incr c_routine] }
+    integer :: x, y
+    integer, parameter :: c = 3
+    integer, external :: g
+
+          ! Can't have nested scopes fun.  (Fortran 'block' construct supported only starting with OpenACC 3.1.)
+
+          ! Don't know how to effect a 'LABEL_DECL' here.
+          ! Don't know how to effect a 'TYPE_DECL' here.
+          ! Don't know how to effect a 'FUNCTION_DECL' here.
+          ! Don't know how to effect a 'RESULT_DECL' here.
+          ! Don't know how to effect a 'VAR_DECL' here.
+          ! (See C/C++ example.)
+
+          a = g (i, j, a, c)
+          x = a
+          !$acc atomic write ! ... to force 'TREE_ADDRESSABLE'.
+          y = a
+  end subroutine f
+    ! { dg-note {variable 'i' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { xfail *-*-* } l_routine$c_routine }
+    ! { dg-note {variable 'j' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { xfail *-*-* } l_routine$c_routine }
+    ! { dg-note {variable 'a' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "TODO" { xfail *-*-* } l_routine$c_routine }
+    ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { xfail *-*-* } l_routine$c_routine }
+    ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "TODO" { xfail *-*-* } l_routine$c_routine }
+end module m