Add diagnostic: "note: OpenACC 'kernels' decomposition: variable '[...]' declared in block made addressable" [PR100280]
Commit Message
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 e5ae22c56152b1a1f4b4e1d7ae04431a9e4710cc
"Add diagnostic: "note: OpenACC 'kernels' decomposition: variable '[...]'
declared in block made addressable" [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
From e5ae22c56152b1a1f4b4e1d7ae04431a9e4710cc Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Tue, 15 Feb 2022 16:54:30 +0100
Subject: [PATCH] Add diagnostic: "note: OpenACC 'kernels' decomposition:
variable '[...]' declared in block made addressable" [PR100280]
Follow-up to commit 9b32c1669aad5459dd053424f9967011348add83
"OpenACC 'kernels' decomposition: Mark variables used in
synthesized data clauses as addressable [PR100280]".
PR middle-end/100280
gcc/
* omp-oacc-kernels-decompose.cc (maybe_build_inner_data_region):
Add diagnostic: "note: OpenACC 'kernels' decomposition: variable
'[...]' declared in block made addressable".
gcc/testsuite/
* c-c++-common/goacc/classify-kernels-unparallelized.c: Add
'--param=openacc-privatization=noisy'.
* c-c++-common/goacc/classify-kernels.c: Likewise.
* c-c++-common/goacc/kernels-decompose-2.c: Adjust.
* 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-oacc-kernels-decompose.cc | 24 ++++++++++++++++++-
.../goacc/classify-kernels-unparallelized.c | 7 ++++++
.../c-c++-common/goacc/classify-kernels.c | 7 ++++++
.../c-c++-common/goacc/kernels-decompose-2.c | 2 ++
.../goacc/kernels-decompose-pr100280-1.c | 1 +
.../goacc/kernels-decompose-pr104061-1-2.c | 1 +
.../goacc/kernels-decompose-pr104061-1-3.c | 1 +
.../goacc/kernels-decompose-pr104061-1-4.c | 1 +
.../goacc/kernels-decompose-pr104132-1.c | 1 +
.../goacc/kernels-decompose-pr104133-1.c | 1 +
.../libgomp.oacc-c-c++-common/f-asyncwait-1.c | 3 +++
.../kernels-decompose-1.c | 1 +
12 files changed, 49 insertions(+), 1 deletion(-)
@@ -845,7 +845,29 @@ maybe_build_inner_data_region (location_t loc, gimple *body,
prev_mapped_var = v;
/* See <https://gcc.gnu.org/PR100280>. */
- TREE_ADDRESSABLE (v) = 1;
+ if (!TREE_ADDRESSABLE (v))
+ {
+ TREE_ADDRESSABLE (v) = 1;
+
+ if (dump_enabled_p ())
+ {
+ 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%> declared in block"
+ " made addressable\n",
+ v);
+#if __GNUC__ >= 10
+# pragma GCC diagnostic pop
+#endif
+ }
+ }
}
}
@@ -9,6 +9,10 @@
{ dg-additional-options "-fdump-tree-parloops1-all" }
{ dg-additional-options "-fdump-tree-oaccloops" } */
+/* { dg-additional-options "--param=openacc-privatization=noisy" }
+ Prune a few: uninteresting:
+ { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
+
/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
aspects of that functionality. */
@@ -25,6 +29,9 @@ 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 {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 } */
/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
for (unsigned int i = 0; i < N; i++)
@@ -9,6 +9,10 @@
{ dg-additional-options "-fdump-tree-parloops1-all" }
{ dg-additional-options "-fdump-tree-oaccloops" } */
+/* { dg-additional-options "--param=openacc-privatization=noisy" }
+ Prune a few: uninteresting:
+ { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
+
/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
aspects of that functionality. */
@@ -21,6 +25,9 @@ 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 {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 } */
/* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
for (unsigned int i = 0; i < N; i++)
@@ -66,12 +66,14 @@ 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 {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 {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 } */
@@ -10,6 +10,7 @@ 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 {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 }
@@ -16,6 +16,7 @@ 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-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 }
@@ -23,6 +23,7 @@ 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: 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 } */
{
@@ -21,6 +21,7 @@ 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: 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 } */
{
@@ -18,6 +18,7 @@ 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 {variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */
{
int k;
@@ -18,6 +18,7 @@ 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 {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 } */
@@ -260,6 +260,7 @@ 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 {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 } */
@@ -268,6 +269,7 @@ 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 {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 } */
@@ -311,6 +313,7 @@ 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 {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 } */
@@ -32,6 +32,7 @@ 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 {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