OpenACC 'kernels' decomposition: Move 'TREE_ADDRESSABLE' setting into OMP lowering...

... in preparation for later changes.  No functional change.

Follow-up to commit 9b32c1669aad5459dd053424f9967011348add83
"OpenACC 'kernels' decomposition: Mark variables used in
synthesized data clauses as addressable [PR100280]".

	PR middle-end/100280
	gcc/
	* tree.h (OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE): New.
	* tree-core.h: Document it.
	* omp-low.cc (scan_sharing_clauses) <OMP_CLAUSE_MAP>: Handle
	'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE'.
	* omp-oacc-kernels-decompose.cc (maybe_build_inner_data_region):
	Set 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' instead of
	'TREE_ADDRESSABLE'.
	gcc/testsuite/
	* c-c++-common/goacc/classify-kernels-unparallelized.c: Adjust.
	* c-c++-common/goacc/classify-kernels.c: Likewise.
	* c-c++-common/goacc/kernels-decompose-2.c: Likewise.
	* c-c++-common/goacc/kernels-decompose-pr100280-1.c: Likewise.
	* c-c++-common/goacc/kernels-decompose-pr104061-1-2.c: Likewise.
	* c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: Likewise.
	* c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: Likewise.
	* c-c++-common/goacc/kernels-decompose-pr104132-1.c: Likewise.
	* c-c++-common/goacc/kernels-decompose-pr104133-1.c: Likewise.
	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c: Adjust.
	* testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
	Likewise.

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 2294456b27d4f8f125fba1692fa45a65640eb829..6654bfd426e092cb41e52bd5056533e5503d4741 100644 (file)
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1495,6 +1495,37 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
          if (ctx->outer)
            scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer);
          decl = OMP_CLAUSE_DECL (c);
+         /* If requested, make 'decl' addressable.  */
+         if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+             && OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (c))
+           {
+             gcc_checking_assert (DECL_P (decl));
+
+             gcc_checking_assert (!TREE_ADDRESSABLE (decl));
+             TREE_ADDRESSABLE (decl) = 1;
+
+             if (dump_enabled_p ())
+               {
+                 location_t loc = OMP_CLAUSE_LOCATION (c);
+                 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,
+                                  "variable %<%T%>"
+                                  " made addressable\n",
+                                  decl);
+#if __GNUC__ >= 10
+# pragma GCC diagnostic pop
+#endif
+               }
+
+             /* Done.  */
+             OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (c) = 0;
+           }
          /* Global variables with "omp declare target" attribute
             don't need to be copied, the receiver side will use them
             directly.  However, global variables with "omp declare target link"

diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc
index 5093386f718682907b71f6af7b05e1f85c656b67..ecbd3071e5dfbb19c3f1ae71f6a5f28324d7c188 100644 (file)
--- a/gcc/omp-oacc-kernels-decompose.cc
+++ b/gcc/omp-oacc-kernels-decompose.cc
@@ -847,7 +847,8 @@ maybe_build_inner_data_region (location_t loc, gimple *body,
          /* See <https://gcc.gnu.org/PR100280>.  */
          if (!TREE_ADDRESSABLE (v))
            {
-             TREE_ADDRESSABLE (v) = 1;
+             /* Request that OMP lowering make 'v' addressable.  */
+             OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (new_clause) = 1;
 
              if (dump_enabled_p ())
                {
@@ -861,7 +862,7 @@ maybe_build_inner_data_region (location_t loc, gimple *body,
                  dump_printf_loc (MSG_NOTE, d_u_loc,
                                   "OpenACC %<kernels%> decomposition:"
                                   " variable %<%T%> declared in block"
-                                  " made addressable\n",
+                                  " requested to be made addressable\n",
                                   v);
 #if __GNUC__ >= 10
 # pragma GCC diagnostic pop

diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c
index 2496462f7770fe91342434907ced3aa5b7b3a6af..61871d118a9800fa35b3a4c46a19028d6fe6edb1 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c
@@ -29,7 +29,8 @@ extern unsigned int f (unsigned int);
 void KERNELS ()
 {
 #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-line l_compute1 } */
-  /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_compute1 } */
+  /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute1 } */
+  /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute1 } */
   /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */
   /* { dg-note {variable '[abc]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */
   /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute1 } */

diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c
index 3ba0411a57ab39b206df2532977802d2c3537c7a..1473337986f9fce695665afc0344fea7b0c0dc91 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c
@@ -25,7 +25,8 @@ extern unsigned int *__restrict c;
 void KERNELS ()
 {
 #pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-line l_compute1 } */
-  /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_compute1 } */
+  /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute1 } */
+  /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute1 } */
   /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */
   /* { dg-note {variable '[abc]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */
   /* { dg-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_compute1 } */

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c
index 5fbd10221a2bc7c0c9d47a281e46549f33e3a749..bf158311daec831a4533678d4e538e0a8ca55c32 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c
@@ -66,14 +66,16 @@ main ()
 
 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
   /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } l_compute$c_compute } */
-  /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block 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 } */
   {
     int i;
   }
 
 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
-  /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block 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 gang loop parallelism} {} { target *-*-* } l_compute$c_compute } */
   /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c
index 91e589777b0dd33acedd8535acc88760339c112d..1c1e22c00ac39ac468653231382e1abcc44b7bde 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100280-1.c
@@ -10,7 +10,8 @@ void
 foo (void) /* { dg-line l_f_1 } */
 {
 #pragma acc kernels /* { dg-line l_k_1 } */
-  /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block made addressable} {} { target *-*-* } l_k_1 } */
+  /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_k_1 } */
+  /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_k_1 } */
   /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_k_1 } */
   /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_k_1 } */
   /* { dg-bogus {note: beginning 'parloops' part in OpenACC 'kernels' region} {TODO location} { xfail *-*-* } l_f_1 }

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c
index 8ca02cb2119a36606ec4df17720b2e68386caaea..336cf2ad4256eabea530fc36fab8e86d42bd50ca 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c
@@ -16,7 +16,8 @@ foo (void)
 {
   /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail *-*-* } .+1 } */
 #pragma acc kernels /* { dg-line l_compute1 } */
-  /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */
+  /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */
+  /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */
   /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_compute1 } */
   {
     /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 }

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c
index 05a196d8f367af37457cf2704cb360613bd7b9c3..f41dda86122f58d2acda87f0618a4cf1f14d9a47 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c
@@ -23,7 +23,8 @@ foo (void)
 {
   /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'.  */
 #pragma acc kernels /* { dg-line l_compute1 } */
-  /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */
+  /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */
+  /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */
   /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_compute1 } */
   /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */
   {

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c
index 07cb592bcf96eb8279b0eb3e49dd46bc7f2a7027..cde95a7b7ac96fe9986d254a66c2933846d6ec63 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c
@@ -21,7 +21,8 @@ foo (void)
 {
   /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'.  */
 #pragma acc kernels /* { dg-line l_compute1 } */
-  /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */
+  /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */
+  /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { xfail *-*-* } l_compute1 } */
   /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } l_compute1 } */
   /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */
   {

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c
index 5b0fe42efbf7887fc0a9848c401767f076d885da..4f38a83bb193e4cefd716c7bb5ad6a986a0f987e 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104132-1.c
@@ -18,7 +18,8 @@ void
 foo (void)
 {
 #pragma acc kernels /* { dg-line l_compute1 } */
-  /* { dg-note {OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {} { target *-*-* } l_compute1 } */
+  /* { dg-note {OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {} { target *-*-* } l_compute1 } */
+  /* { dg-note {variable 'k' made addressable} {} { target *-*-* } l_compute1 } */
   /* { dg-note {variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */
   {
     int k;

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c
index 4536b5c179ebe4c242f60fe201ebe126faecf4e2..0499665777dcb7c005e81b34f2e21a57fa954e1a 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104133-1.c
@@ -18,7 +18,8 @@ void
 foo (void)
 {
 #pragma acc kernels /* { dg-line l_compute1 } */
-  /* { dg-note {OpenACC 'kernels' decomposition: variable 'k' declared in block made addressable} {} { target *-*-* } l_compute1 } */
+  /* { dg-note {OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {} { target *-*-* } l_compute1 } */
+  /* { dg-note {variable 'k' made addressable} {} { target *-*-* } l_compute1 } */
   /* { dg-note {variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */
   /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */
   /* { dg-note {variable 'arr_0\.1' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */

diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 01a1ce499daad90c06c5a11d1e83d30592d591f8..4530bd8c2c7633bed2223c495210e552ff329e0f 100644 (file)
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -1155,6 +1155,9 @@ struct GTY(()) tree_base {
        PREDICT_EXPR_OUTCOME in
           PREDICT_EXPR
 
+       OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE in
+          OMP_CLAUSE
+
    static_flag:
 
        TREE_STATIC in

diff --git a/gcc/tree.h b/gcc/tree.h
index 36ceed57064048bf252d6724bfc41cf8c9d2b73c..da6f3b38ba556c315c9abb145e3e74090582ce09 100644 (file)
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1699,6 +1699,11 @@ class auto_suppress_location_wrappers
 #define OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P(NODE) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.deprecated_flag)
 
+/* Flag that 'OMP_CLAUSE_DECL (NODE)' is to be made addressable during OMP
+   lowering.  */
+#define OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE(NODE) \
+  (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.addressable_flag)
+
 /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present'
    clause.  */
 #define OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT(NODE) \

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c
index d360aad4736d5652449c99b28e5da3300c54d57e..9f7b8b4af2fbe99db8f2b5544d295498cde2f20f 100644 (file)
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c
@@ -260,7 +260,8 @@ main (void)
 #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N])
   {
 #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */
-    /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block 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 { ! __OPTIMIZE__ } } l_compute$c_compute }
        { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */
@@ -269,7 +270,8 @@ main (void)
       b[i] = (a[i] * a[i] * a[i]) / a[i];
 
 #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */
-    /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block 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 { ! __OPTIMIZE__ } } l_compute$c_compute }
        { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */
@@ -313,7 +315,8 @@ main (void)
 #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N])
   {
 #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */
-    /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block 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 { ! __OPTIMIZE__ } } l_compute$c_compute }
        { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */

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 b99497ea6c59bbfdfcb1cabc0d0269d87cef5e07..85c39871f94551c08b1812cb24cc0eb7743bd2db 100644 (file)
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c
@@ -32,7 +32,8 @@ int main()
   {
     /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
     int c = 234;
-    /* { dg-note {OpenACC 'kernels' decomposition: variable 'c' declared in block made addressable} "" { target *-*-* } l_compute$c_compute } */
+    /* { dg-note {OpenACC 'kernels' decomposition: variable 'c' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */
+    /* { dg-note {variable 'c' made addressable} {} { target *-*-* } l_compute$c_compute } */
     /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */
 
 #pragma acc loop independent gang /* { dg-line l_loop_i[incr c_loop_i] } */