OpenACC 'kernels' decomposition: Mark variables used in 'present' clauses as addressable [PR100280, PR104086]

... like in recent commit 9b32c1669aad5459dd053424f9967011348add83
"OpenACC 'kernels' decomposition: Mark variables used in synthesized
data clauses as addressable [PR100280]".  Otherwise, we may run into
'gcc/omp-low.cc:lower_omp_target':

    13125                       else if (is_gimple_reg (var))
    13126                         {
    13127                           gcc_assert (offloaded);

	PR middle-end/100280
	PR middle-end/104086
	gcc/
	* omp-oacc-kernels-decompose.cc (omp_oacc_kernels_decompose_1):
	Mark variables used in 'present' clauses as addressable.
	* omp-low.cc (scan_sharing_clauses) <OMP_CLAUSE_MAP>: Gracefully
	handle duplicate 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE'.
	gcc/testsuite/
	* c-c++-common/goacc/kernels-decompose-pr104086-1.c: Adjust,
	extend.
	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c:
	Merge this...
	* testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c:
	..., and this...
	* testsuite/libgomp.oacc-c-c++-common/declare-vla.c: ... into
	this, and adjust.
	* testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
	Extend.

3 files changed, 80 insertions, 16 deletions

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index d932d74..cfc63d6 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1501,11 +1501,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	    {
 	      gcc_checking_assert (DECL_P (decl));
 
-	      gcc_checking_assert (!TREE_ADDRESSABLE (decl));
-	      if (!make_addressable_vars)
-		make_addressable_vars = BITMAP_ALLOC (NULL);
-	      bitmap_set_bit (make_addressable_vars, DECL_UID (decl));
-	      TREE_ADDRESSABLE (decl) = 1;
+	      bool decl_addressable = TREE_ADDRESSABLE (decl);
+	      if (!decl_addressable)
+		{
+		  if (!make_addressable_vars)
+		    make_addressable_vars = BITMAP_ALLOC (NULL);
+		  bitmap_set_bit (make_addressable_vars, DECL_UID (decl));
+		  TREE_ADDRESSABLE (decl) = 1;
+		}
 
 	      if (dump_enabled_p ())
 		{
@@ -1517,10 +1520,16 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 # pragma GCC diagnostic push
 # pragma GCC diagnostic ignored "-Wformat"
 #endif
-		  dump_printf_loc (MSG_NOTE, d_u_loc,
-				   "variable %<%T%>"
-				   " made addressable\n",
-				   decl);
+		  if (!decl_addressable)
+		    dump_printf_loc (MSG_NOTE, d_u_loc,
+				     "variable %<%T%>"
+				     " made addressable\n",
+				     decl);
+		  else
+		    dump_printf_loc (MSG_NOTE, d_u_loc,
+				     "variable %<%T%>"
+				     " already made addressable\n",
+				     decl);
 #if __GNUC__ >= 10
 # pragma GCC diagnostic pop
 #endif
diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc
index ecbd307..40b0453 100644
--- a/gcc/omp-oacc-kernels-decompose.cc
+++ b/gcc/omp-oacc-kernels-decompose.cc
@@ -1468,6 +1468,38 @@ omp_oacc_kernels_decompose_1 (gimple *kernels_stmt)
 		  /* Now that this data is mapped, turn the data clause on the
 		     inner OpenACC 'kernels' into a 'present' clause.  */
 		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_PRESENT);
+
+		  /* See <https://gcc.gnu.org/PR100280>,
+		     <https://gcc.gnu.org/PR104086>.  */
+		  if (DECL_P (decl)
+		      && !TREE_ADDRESSABLE (decl))
+		    {
+		      /* Request that OMP lowering make 'decl' addressable.  */
+		      OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (new_clause) = 1;
+
+		      if (dump_enabled_p ())
+			{
+			  location_t loc = OMP_CLAUSE_LOCATION (new_clause);
+			  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
+			    (MSG_NOTE, d_u_loc,
+			     "OpenACC %<kernels%> decomposition:"
+			     " variable %<%T%> in %qs clause"
+			     " requested to be made addressable\n",
+			     decl,
+			     user_omp_clause_code_name (new_clause, true));
+#if __GNUC__ >= 10
+# pragma GCC diagnostic pop
+#endif
+			}
+		    }
 		}
 	      break;
 
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104086-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104086-1.c
index eab10cf..83fb75e 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104086-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104086-1.c
@@ -1,8 +1,5 @@
-/* Reduced from 'libgomp.oacc-c-c++-common/declare-vla.c'.  */
-
-/* { dg-additional-options "-fchecking" }
-   { dg-ice TODO }
-   { dg-prune-output {during GIMPLE pass: omplower} } */
+/* Reduced from 'libgomp.oacc-c-c++-common/declare-vla.c', and then
+   extended.  */
 
 /* { dg-additional-options "--param openacc-kernels=decompose" } */
 
@@ -14,12 +11,38 @@ void
 foo (void)
 {
 #pragma acc data /* { dg-line l_data1 } */
-  /* { dg-bogus {note: variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO 'data'} { xfail *-*-* } l_data1 } */
+  /* { dg-bogus {note: variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {TODO 'data'} { xfail *-*-* } l_data1 } */
   {
     int i;
 
-#pragma acc kernels
+#pragma acc kernels /* { dg-line l_compute1 } */
+    /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 }
+       { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute1 } */
     /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
     i = 0;
+
+#pragma acc kernels /* { dg-line l_compute2 } */
+    /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute2 }
+       { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute2 } */
+    /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+    i = -1;
   }
 }
+
+void
+foo2 (void)
+{
+  int i[1];
+
+#pragma acc kernels /* { dg-line l2_compute1 } */
+  /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l2_compute1 }
+     { dg-note {variable 'i' made addressable} {} { target *-*-* } l2_compute1 } */
+  /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+  i[0] = 0;
+
+#pragma acc kernels /* { dg-line l2_compute2 } */
+  /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l2_compute2 }
+     { dg-note {variable 'i' already made addressable} {} { target *-*-* } l2_compute2 } */
+  /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+  i[0] = -1;
+}