diff options
author | Thomas Schwinge <thomas@codesourcery.com> | 2022-02-15 16:54:30 +0100 |
---|---|---|
committer | Thomas Schwinge <thomas@codesourcery.com> | 2022-03-04 14:21:00 +0100 |
commit | e5ae22c56152b1a1f4b4e1d7ae04431a9e4710cc (patch) | |
tree | 5ad220245bc1531acbe0feafc450de824fa8dfeb | |
parent | c14ea6a72fb1ae66e3d32ac8329558497c6e4403 (diff) | |
download | gcc-e5ae22c56152b1a1f4b4e1d7ae04431a9e4710cc.zip gcc-e5ae22c56152b1a1f4b4e1d7ae04431a9e4710cc.tar.gz gcc-e5ae22c56152b1a1f4b4e1d7ae04431a9e4710cc.tar.bz2 |
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.
12 files changed, 49 insertions, 1 deletions
diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc index 98eafdb..5093386 100644 --- a/gcc/omp-oacc-kernels-decompose.cc +++ b/gcc/omp-oacc-kernels-decompose.cc @@ -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 + } + } } } 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 4ee8e9d..2496462 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c @@ -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++) diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c index 74acca8..3ba0411 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c @@ -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++) 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 64ce894..5fbd102 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c @@ -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 } */ 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 b497af2..91e5897 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,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 } 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 ea58ea4..8ca02cb 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,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 } 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 ee3fcc7..05a196d 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,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 } */ { 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 c140e61..07cb592 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,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 } */ { 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 4fbfdd8..5b0fe42 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,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; 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 72dde34..4536b5c 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,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 } */ 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 ef7735b..d360aad 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,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 } */ 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 57e75f6..b99497e 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,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] } */ |