From: Thomas Schwinge <thomas@codesourcery.com>
To: <gcc-patches@gcc.gnu.org>
Cc: <fortran@gcc.gnu.org>, Jakub Jelinek <jakub@redhat.com>,
Julian Brown <julian@codesourcery.com>, <asolokha@gmx.com>
Subject: OpenACC 'kernels' decomposition: Move 'TREE_ADDRESSABLE' setting into OMP lowering [PR100280]
Date: Fri, 4 Mar 2022 14:41:47 +0100 [thread overview]
Message-ID: <878rtpomw4.fsf@euler.schwinge.homeip.net> (raw)
In-Reply-To: <87iluovu47.fsf@euler.schwinge.homeip.net>
[-- Attachment #1: Type: text/plain, Size: 2589 bytes --]
Hi!
On 2022-01-13T10:54:16+0100, I wrote:
> On 2019-05-08T14:51:57+0100, Julian Brown <julian@codesourcery.com> wrote:
>> - The "addressable" bit is set during the kernels conversion pass for
>> variables that have "create" (alloc) clauses created for them in the
>> synthesised outer data region (instead of in the front-end, etc.,
>> where it can't be done accurately). Such variables actually have
>> their address taken during transformations made in a later pass
>> (omp-low, I think), but there's a phase-ordering problem that means
>> the flag should be set earlier.
>
> The actual issue is a bit different, but yes, there is a problem.
> The related ICE has also been reported as <https://gcc.gnu.org/PR100280>
> "ICE in lower_omp_target, at omp-low.c:12287". (And I'm confused why we
> didn't run into that with the OpenACC 'kernels' decomposition
> originally.) I've pushed to master branch
> commit 9b32c1669aad5459dd053424f9967011348add83
> "OpenACC 'kernels' decomposition: Mark variables used in synthesized data
> clauses as addressable [PR100280]", see attached.
> --- a/gcc/omp-oacc-kernels-decompose.cc
> +++ b/gcc/omp-oacc-kernels-decompose.cc
> @@ -793,7 +793,8 @@ make_data_region_try_statement (location_t loc, gimple *body)
>
> /* If INNER_BIND_VARS holds variables, build an OpenACC data region with
> location LOC containing BODY and having 'create (var)' clauses for each
> - variable. If INNER_CLEANUP is present, add a try-finally statement with
> + variable (as a side effect, such variables also get TREE_ADDRESSABLE set).
> + If INNER_CLEANUP is present, add a try-finally statement with
> this cleanup code in the finally block. Return the new data region, or
> the original BODY if no data region was needed. */
>
> @@ -842,6 +843,9 @@ maybe_build_inner_data_region (location_t loc, gimple *body,
> inner_data_clauses = new_clause;
>
> prev_mapped_var = v;
> +
> + /* See <https://gcc.gnu.org/PR100280>. */
> + TREE_ADDRESSABLE (v) = 1;
> }
> }
Pushed to master branch commit de6e81ea961219d0726db67776d11ce75a4cae1b
"OpenACC 'kernels' decomposition: Move 'TREE_ADDRESSABLE' setting into
OMP lowering [PR100280]", see attached.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-OpenACC-kernels-decomposition-Move-TREE_ADDRESSABLE-.patch --]
[-- Type: text/x-diff, Size: 21643 bytes --]
From de6e81ea961219d0726db67776d11ce75a4cae1b Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Tue, 15 Feb 2022 23:03:49 +0100
Subject: [PATCH] OpenACC 'kernels' decomposition: Move 'TREE_ADDRESSABLE'
setting into OMP lowering [PR100280]
... 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.
---
gcc/omp-low.cc | 31 +++++++++++++++++++
gcc/omp-oacc-kernels-decompose.cc | 5 +--
.../goacc/classify-kernels-unparallelized.c | 3 +-
.../c-c++-common/goacc/classify-kernels.c | 3 +-
.../c-c++-common/goacc/kernels-decompose-2.c | 6 ++--
.../goacc/kernels-decompose-pr100280-1.c | 3 +-
.../goacc/kernels-decompose-pr104061-1-2.c | 3 +-
.../goacc/kernels-decompose-pr104061-1-3.c | 3 +-
.../goacc/kernels-decompose-pr104061-1-4.c | 3 +-
.../goacc/kernels-decompose-pr104132-1.c | 3 +-
.../goacc/kernels-decompose-pr104133-1.c | 3 +-
gcc/tree-core.h | 3 ++
gcc/tree.h | 5 +++
.../libgomp.oacc-c-c++-common/f-asyncwait-1.c | 9 ++++--
.../kernels-decompose-1.c | 3 +-
15 files changed, 70 insertions(+), 16 deletions(-)
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 2294456b27d..6654bfd426e 100644
--- 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 5093386f718..ecbd3071e5d 100644
--- 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 2496462f777..61871d118a9 100644
--- 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 3ba0411a57a..1473337986f 100644
--- 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 5fbd10221a2..bf158311dae 100644
--- 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 91e589777b0..1c1e22c00ac 100644
--- 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 8ca02cb2119..336cf2ad425 100644
--- 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 05a196d8f36..f41dda86122 100644
--- 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 07cb592bcf9..cde95a7b7ac 100644
--- 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 5b0fe42efbf..4f38a83bb19 100644
--- 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 4536b5c179e..0499665777d 100644
--- 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 01a1ce499da..4530bd8c2c7 100644
--- 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 36ceed57064..da6f3b38ba5 100644
--- 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 d360aad4736..9f7b8b4af2f 100644
--- 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 b99497ea6c5..85c39871f94 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
@@ -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] } */
--
2.34.1
prev parent reply other threads:[~2022-03-04 13:41 UTC|newest]
Thread overview: 10+ messages / expand[flat|nested] mbox.gz Atom feed top
[not found] <20190508145157.08beb4df@squid.athome>
2022-01-13 9:54 ` OpenACC 'kernels' decomposition: Mark variables used in synthesized data clauses as addressable [PR100280] Thomas Schwinge
2022-03-01 16:46 ` Thomas Schwinge
2022-03-02 13:30 ` Jakub Jelinek
2022-03-04 13:34 ` Add 'c-c++-common/goacc/kernels-decompose-pr104132-1.c' [PR104132] Thomas Schwinge
2022-03-04 13:34 ` Add 'c-c++-common/goacc/kernels-decompose-pr104133-1.c' [PR104133] Thomas Schwinge
2022-03-04 13:46 ` OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs [PR100280, PR104132, PR104133] Thomas Schwinge
2022-03-04 14:57 ` Test 'libgomp.oacc-*/kernels-private-vars-*' with '--param=openacc-kernels=decompose' [PR104784] Thomas Schwinge
2022-03-04 19:52 ` Thomas Schwinge
2022-03-04 13:37 ` Add diagnostic: "note: OpenACC 'kernels' decomposition: variable '[...]' declared in block made addressable" [PR100280] Thomas Schwinge
2022-03-04 13:41 ` Thomas Schwinge [this message]
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=878rtpomw4.fsf@euler.schwinge.homeip.net \
--to=thomas@codesourcery.com \
--cc=asolokha@gmx.com \
--cc=fortran@gcc.gnu.org \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
--cc=julian@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).