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.

4 files changed, 88 insertions, 67 deletions

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c
deleted file mode 100644
index 3e5b6bab233..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c
+++ /dev/null
@@ -1,22 +0,0 @@
-/* { dg-additional-options "--param=openacc-kernels=decompose" } */
-/* ICE similar to PR100280, but not the same.
-   { dg-ice "TODO" }
-   TODO { dg-prune-output "during GIMPLE pass: omplower" }
-   TODO { dg-do link } */
-
-/* { dg-additional-options "-fopt-info-omp-all" }
-   { dg-additional-options "-foffload=-fopt-info-all-omp" } */
-
-/* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
-   Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
-   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
-
-#undef KERNELS_DECOMPOSE_ICE_HACK
-#include "declare-vla.c"
-
-/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } 27 } */
-
-/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } 61 } */
-
-/* { dg-bogus {note: variable [^\n\r]+ candidate for adjusting OpenACC privatization level} {TODO 'data'} { xfail *-*-* } 42 } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c
deleted file mode 100644
index 142aceec9cd..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c
+++ /dev/null
@@ -1,29 +0,0 @@
-/* { dg-additional-options "--param=openacc-kernels=decompose" } */
-
-/* See also 'declare-vla-kernels-decompose-ice-1.c'.  */
-
-/* { dg-additional-options "-fopt-info-omp-all" }
-   { dg-additional-options "-foffload=-fopt-info-all-omp" } */
-
-/* { dg-additional-options "--param=openacc-privatization=noisy" }
-   { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
-   Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
-   { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
-
-#define KERNELS_DECOMPOSE_ICE_HACK
-#include "declare-vla.c"
-
-/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } 27 } */
-
-/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } 61 } */
-
-/* { dg-bogus {note: variable [^\n\r]+ candidate for adjusting OpenACC privatization level} {TODO 'data'} { xfail *-*-* } 42 } */
-
-/* { dg-note {variable 'i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } 58 }
-   { dg-note {variable 'N\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } 58 } */
-
-/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } 24 }
-   { dg-optimized {assigned OpenACC gang loop parallelism} {} { target { __OPTIMIZE__ } } 24 } */
-
-/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } 58 }
-   { dg-optimized {assigned OpenACC gang loop parallelism} {} { target { __OPTIMIZE__ } } 58 } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
index 4ce2e6d1f18..f6fc3ffefa4 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
@@ -1,5 +1,7 @@
 /* Verify OpenACC 'declare' with VLAs.  */
 
+/* { dg-additional-options "--param=openacc-kernels=decompose" } */
+
 /* { dg-additional-options "-fopt-info-omp-all" }
    { dg-additional-options "-foffload=-fopt-info-all-omp" } */
 
@@ -8,6 +10,15 @@
    Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
    { 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 [...]",
+   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".  */
+
+
 #include <assert.h>
 
 
@@ -21,9 +32,10 @@ f (void)
   for (i = 0; i < N; i++)
     A[i] = -i;
 
-#pragma acc kernels
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } .-1 }
-     { dg-optimized {assigned OpenACC gang loop parallelism} {} { target { __OPTIMIZE__ } } .-2 } */
+#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
+  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } l_compute$c_compute }
+     { dg-optimized {assigned OpenACC gang loop parallelism} {} { target __OPTIMIZE__ } l_compute$c_compute } */
+  /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
   for (i = 0; i < N; i++)
     A[i] = i;
 
@@ -49,15 +61,14 @@ f_data (void)
     for (i = 0; i < N; i++)
       A[i] = -i;
 
-    /* See 'declare-vla-kernels-decompose.c'.  */
-#ifdef KERNELS_DECOMPOSE_ICE_HACK
-    (volatile int *) &i;
-    (volatile int *) &N;
-#endif
-
-# pragma acc kernels
-  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } .-1 }
-     { dg-optimized {assigned OpenACC gang loop parallelism} {} { target { __OPTIMIZE__ } } .-2 } */
+# pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
+    /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+       { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
+    /* { dg-note {OpenACC 'kernels' decomposition: variable 'N' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+       { dg-note {variable 'N' made addressable} {} { target *-*-* } l_compute$c_compute } */
+    /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target { ! __OPTIMIZE__ } } l_compute$c_compute }
+       { dg-optimized {assigned OpenACC gang loop parallelism} {} { target __OPTIMIZE__ } l_compute$c_compute } */
+    /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
     for (i = 0; i < N; i++)
       A[i] = i;
 
@@ -78,6 +89,3 @@ main ()
 
   return 0;
 }
-
-
-/* { dg-note dummy "" { target n-on-e } } to disable 'prune_notes'.  */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c
index 40786c750d1..eb424776b6b 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c
@@ -24,7 +24,9 @@
 static int g1;
 static int g2;
 
-int main()
+/* PR100280, etc. */
+
+static void f1 ()
 {
   int a = 0;
   /*TODO Without making 'a' addressable, for GCN offloading we will not see the expected value copied out.  (But it does work for nvptx offloading, strange...)  */
@@ -153,5 +155,67 @@ int main()
   assert (g2 == N * (N + 1) / 2);
   assert (f1 == 2432902008176640000ULL);
 
+#undef N
+}
+
+
+/* PR104086 */
+
+static void f2 ()
+{
+#pragma acc data
+  /* { dg-bogus {note: variable [^\n\r]+ candidate for adjusting OpenACC privatization level} {TODO 'data'} { xfail *-*-* } .-1 } */
+  {
+    int i;
+
+#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
+    /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+       { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
+    /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+    i = 1;
+
+    assert (i == 1);
+
+#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
+    /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+       { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+    /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+    i = -1;
+
+    assert (i == -1);
+  }
+
+
+  int ia[1];
+
+#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
+  /* { dg-note {OpenACC 'kernels' decomposition: variable 'ia' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+     { dg-note {variable 'ia' made addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+  ia[0] = -2;
+
+  assert (ia[0] == -2);
+
+#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
+  /* { dg-note {OpenACC 'kernels' decomposition: variable 'ia' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+     { dg-note {variable 'ia' already made addressable} {} { target *-*-* } l_compute$c_compute } */
+  /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
+     { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute }
+     { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */
+  /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
+  /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+  for (int i = 0; i < 100; ++i)
+    ++ia[0];
+
+  assert (ia[0] == -2 + 100);
+}
+
+
+int main()
+{
+  f1 ();
+
+  f2 ();
+
   return 0;
 }