diff options
22 files changed, 120 insertions, 45 deletions
diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc index 40b04539894..4386787ba3c 100644 --- a/gcc/omp-oacc-kernels-decompose.cc +++ b/gcc/omp-oacc-kernels-decompose.cc @@ -1505,7 +1505,6 @@ omp_oacc_kernels_decompose_1 (gimple *kernels_stmt) case GOMP_MAP_POINTER: case GOMP_MAP_TO_PSET: - case GOMP_MAP_FORCE_TOFROM: case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FIRSTPRIVATE_REFERENCE: /* ??? Copying these map kinds leads to internal compiler 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 bf158311dae..3ce9490f02f 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-2.c @@ -46,7 +46,13 @@ main () int a[N], b[N], c[N]; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'x\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'z' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'z' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'y' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'y' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'x' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'x' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'x\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ { /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ x = 0; @@ -58,6 +64,8 @@ main () { 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-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ for (i = 0; i < N; i++) @@ -66,16 +74,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 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 '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 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 '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 } */ @@ -90,7 +98,9 @@ main () for (int i = 0; i < N; i++) b[i] = a[N - i - 1]; -#pragma acc kernels +#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'z' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'z' already made addressable} {} { target *-*-* } l_compute$c_compute } */ { #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */ @@ -129,6 +139,8 @@ main () } #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'y' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'y' already made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ /*TODO What does this mean? TODO { dg-optimized "assigned OpenACC worker vector loop parallelism" "" { target *-*-* } l_compute$c_compute } */ @@ -166,6 +178,8 @@ main () } #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'y' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'y' already made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'kernels'" { xfail *-*-* } l_compute$c_compute } */ { y = f_g (a[5]); /* { dg-line l_part[incr c_part] } */ @@ -182,7 +196,11 @@ main () b[j] = y + f_w (c[j]); /* { dg-optimized "assigned OpenACC worker vector loop parallelism" } */ } -#pragma acc kernels +#pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'z' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'z' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'y' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'y' already made addressable} {} { target *-*-* } l_compute$c_compute } */ { /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ y = 3; diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-1.c index b1b094f9d4f..57cb1a8cb87 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-1.c @@ -13,6 +13,8 @@ void foo (void) { #pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'p' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } + { dg-note {variable 'p' made addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'c\.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-pr100400-1-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-2.c index 0bc4844560a..a643f109bf1 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-2.c @@ -18,6 +18,8 @@ foo (void) { /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail *-*-* } .+1 } */ #pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'p' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } + { dg-note {variable 'p' made addressable} {} { target *-*-* xfail c++ } l_compute1 } */ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { xfail *-*-* } l_compute1 } */ /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_compute1 } */ { diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-3.c index 6c2cbf88305..9779f1036f6 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-3.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-3.c @@ -20,6 +20,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-note {OpenACC 'kernels' decomposition: variable 'p' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } + { dg-note {variable 'p' made addressable} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute1 } */ /* { dg-note {variable 'c\.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-pr100400-1-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-4.c index fc66a3e7feb..2feabf3c7d6 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-4.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-4.c @@ -18,6 +18,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-note {OpenACC 'kernels' decomposition: variable 'p' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } + { dg-note {variable 'p' made addressable} {} { target *-*-* xfail c++ } l_compute1 } */ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { xfail c++ } l_compute1 } */ /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } l_compute1 } */ { diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-1.c index 35ae81f11a2..aa0fca7b6ed 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-1.c @@ -15,6 +15,8 @@ void foo (void) { #pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'arr_0' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } + { dg-note {variable 'arr_0' made addressable} {} { 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 } */ { int k; 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 336cf2ad425..4d7cbb04f4b 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,8 @@ foo (void) { /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail *-*-* } .+1 } */ #pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'arr_0' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } + { dg-note {variable 'arr_0' made addressable} {} { target *-*-* } 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 } */ 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 28d26e566f6..70c2ac5b531 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 @@ -17,6 +17,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-note {OpenACC 'kernels' decomposition: variable 'arr_0' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } + { dg-note {variable 'arr_0' made addressable} {} { target *-*-* } 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 } */ 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 4d125b5db87..d1cc1a97c9f 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 @@ -17,6 +17,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-note {OpenACC 'kernels' decomposition: variable 'arr_0' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } + { dg-note {variable 'arr_0' made addressable} {} { target *-*-* } 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 } */ 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 36a43ca6d1a..2a663e0ea19 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 @@ -12,6 +12,8 @@ void foo (void) { #pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'arr_0' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } + { dg-note {variable 'arr_0' 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 } */ 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 d9da9dae14c..2724e22a550 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 @@ -12,6 +12,8 @@ void foo (void) { #pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'arr_0' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } + { dg-note {variable 'arr_0' 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 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104774-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104774-1.c index 42faa48f991..3ef0c897bcd 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104774-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104774-1.c @@ -12,6 +12,8 @@ void foo (void) { #pragma acc kernels /* { dg-line l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'arr_0' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } + { dg-note {variable 'arr_0' 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 } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 index 3fe9b34c9c8..2ed6cdb6115 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 @@ -21,6 +21,8 @@ program main call setup(a, b) !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-line l_compute1 } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } */ + ! { dg-note {variable 'i' made 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 } do i = 0, n - 1 diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 index 59990135c75..f6228b97944 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-2.f95 @@ -40,7 +40,15 @@ program main integer, parameter :: N = 10 integer :: a(N), b(N), c(N) - !$acc kernels + !$acc kernels ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'z' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'z' made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'y_l' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'y_l' made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'y' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'y' made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'x' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'x' made addressable} {} { target *-*-* } l_compute$c_compute } ! { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } x = 0 y = 0 @@ -51,6 +59,8 @@ program main !$acc end kernels !$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-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_compute$c_compute } ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } do i = 1, N @@ -66,7 +76,9 @@ program main b(i) = a(N - i + 1) end do - !$acc kernels + !$acc kernels ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'z' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'z' already made addressable} {} { target *-*-* } l_compute$c_compute } !$acc loop ! { dg-line l_loop_i[incr c_loop_i] } ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } @@ -104,6 +116,8 @@ program main !$acc end kernels !$acc kernels ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'y' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'y' already made addressable} {} { target *-*-* } l_compute$c_compute } !TODO What does this mean? !TODO { dg-optimized "assigned OpenACC worker vector loop parallelism" "" { target *-*-* } l_compute$c_compute } !$acc loop independent ! { dg-line l_loop_i[incr c_loop_i] } @@ -141,6 +155,8 @@ program main !$acc end kernels !$acc kernels ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'y' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'y' already made addressable} {} { target *-*-* } l_compute$c_compute } ! { dg-bogus "\[Ww\]arning: region contains gang partitioned code but is not gang partitioned" "TODO 'kernels'" { xfail *-*-* } l_compute$c_compute } y = f_g (a(5)) ! { dg-line l_part[incr c_part] } !TODO If such a construct is placed in its own part (like it is, here), can't this actually use gang paralelism, instead of "gang-single"? @@ -156,7 +172,11 @@ program main end do !$acc end kernels - !$acc kernels + !$acc kernels ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'z' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'z' already made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'y' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'y' already made addressable} {} { target *-*-* } l_compute$c_compute } ! { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } y = 3 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 f6fc3ffefa4..cf423d60327 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c @@ -33,6 +33,10 @@ f (void) A[i] = -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 {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 } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c index fed65c8dccc..9a504384c69 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c @@ -65,8 +65,6 @@ int test_parallel () int test_kernels () { int val = 2; - /*TODO <https://gcc.gnu.org/PR104892> */ - (volatile int *) &val; int ary[32]; int ondev = 0; @@ -75,8 +73,9 @@ int test_kernels () /* val defaults to copy, ary defaults to copy. */ #pragma acc kernels copy(ondev) /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'val' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'val' made addressable} {} { target *-*-* } l_compute$c_compute } */ /* { dg-note {variable 'ondev\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ - /* { dg-note {variable 'val\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ { /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ ondev = acc_on_device (acc_device_not_host); 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 3db59e8a75c..763f697deec 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 @@ -29,15 +29,19 @@ static int g2; static void f1 () { int a = 0; - /*TODO <https://gcc.gnu.org/PR104892> */ - (volatile int *) &a; #define N 123 int b[N] = { 0 }; unsigned long long f1; - /*TODO <https://gcc.gnu.org/PR104892> */ - (volatile void *) &f1; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'f1' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'f1' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'a' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'a' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'g2' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'g2' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'g1' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'g1' made addressable} {} { target *-*-* } l_compute$c_compute } */ { /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ int c = 234; @@ -84,14 +88,12 @@ static void f1 () /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ f1 = 1; - /* { dg-note {variable 'f1\.1' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ #pragma acc loop /* { dg-line l_loop_c[incr c_loop_c] } */ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_c$c_loop_c } */ /* { dg-note {variable 'c' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_c$c_loop_c } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */ for (c = 20; c > 0; --c) f1 *= c; - /* { dg-note {variable 'f1\.2' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ { /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ @@ -109,7 +111,6 @@ static void f1 () { /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ if (f2 != f1) - /* { dg-note {variable 'f1\.3' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target { ! __OPTIMIZE__ } } l_compute$c_compute } */ __builtin_abort (); /* As this is still in the preceding 'parloops' part: @@ -133,7 +134,6 @@ static void f1 () /* As this is still in the preceding 'parloops' part: { dg-bogus {note: beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ if (f2 != f1) - /* { dg-note {variable 'f1\.4' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target { ! __OPTIMIZE__ } } l_compute$c_compute } */ __builtin_abort (); } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-reduction-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-reduction-1.c index e7b2817a391..3da1a495411 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-reduction-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-reduction-1.c @@ -19,11 +19,10 @@ int main () { int i, red = 0; - /*TODO <https://gcc.gnu.org/PR104892> */ - (volatile int *) &red; #pragma acc kernels /* { dg-line l_compute1 } */ - /* { dg-note {variable 'red\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'red' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } + { dg-note {variable 'red' made addressable} {} { target *-*-* } l_compute1 } */ { #pragma acc loop reduction (+:red) /* { dg-line l_loop_i1 } */ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i1 } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c index 75e8cb510cc..b1cfe37df8a 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c @@ -642,14 +642,21 @@ int main () kernels. */ { int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max; - /*TODO <https://gcc.gnu.org/PR104892> */ - (volatile int *) &gangs_min, &gangs_max, &workers_min, &workers_max, &vectors_min, &vectors_max; gangs_min = workers_min = vectors_min = INT_MAX; gangs_max = workers_max = vectors_max = INT_MIN; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'gangs_max\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ - /* { dg-note {variable 'workers_max\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ - /* { dg-note {variable 'vectors_max\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'vectors_max' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'vectors_max' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'vectors_min' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'vectors_min' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'workers_max' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'workers_max' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'workers_min' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'workers_min' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'gangs_max' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'gangs_max' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'gangs_min' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'gangs_min' made addressable} {} { target *-*-* } l_compute$c_compute } */ { #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */ \ reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) @@ -682,17 +689,24 @@ int main () #define WORKERS 5 #define VECTORS 13 int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max; - /*TODO <https://gcc.gnu.org/PR104892> */ - (volatile int *) &gangs_min, &gangs_max, &workers_min, &workers_max, &vectors_min, &vectors_max; gangs_min = workers_min = vectors_min = INT_MAX; gangs_max = workers_max = vectors_max = INT_MIN; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ \ num_gangs (gangs) \ num_workers (WORKERS) \ vector_length (VECTORS) - /* { dg-note {variable 'gangs_max\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ - /* { dg-note {variable 'workers_max\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ - /* { dg-note {variable 'vectors_max\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'vectors_max' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'vectors_max' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'vectors_min' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'vectors_min' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'workers_max' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'workers_max' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'workers_min' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'workers_min' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'gangs_max' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'gangs_max' made addressable} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {OpenACC 'kernels' decomposition: variable 'gangs_min' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + { dg-note {variable 'gangs_min' made addressable} {} { target *-*-* } l_compute$c_compute } */ { #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */ \ reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) diff --git a/libgomp/testsuite/libgomp.oacc-fortran/asyncwait-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/asyncwait-1.f90 index 9440cd7f1b5..6e4e447699d 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/asyncwait-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/asyncwait-1.f90 @@ -219,6 +219,8 @@ program asyncwait !$acc data copy (a(1:N)) copy (b(1:N)) copy (c(1:N)) copy (d(1:N)) !$acc kernels async (1) ! { 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-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_compute$c_compute } ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } do i = 1, N @@ -227,6 +229,8 @@ program asyncwait !$acc end kernels !$acc kernels async (1) ! { 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-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_compute$c_compute } ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } do i = 1, N @@ -263,6 +267,8 @@ program asyncwait !$acc data copy (a(1:N), b(1:N), c(1:N), d(1:N), e(1:N)) !$acc kernels async (1) ! { 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-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_compute$c_compute } ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } do i = 1, N diff --git a/libgomp/testsuite/libgomp.oacc-fortran/kernels-reduction-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/kernels-reduction-1.f90 index 89bae49c94c..0688dd8ae0d 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/kernels-reduction-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/kernels-reduction-1.f90 @@ -15,13 +15,12 @@ program reduction integer, parameter :: n = 20 integer :: i, red - !TODO <https://gcc.gnu.org/PR104892> - call make_addressable (red) red = 0 !$acc kernels ! { dg-line l_compute1 } */ - ! { dg-note {variable 'red\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'red' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute1 } + ! { dg-note {variable 'red' made addressable} {} { target *-*-* } l_compute1 } !$acc loop reduction (+:red) ! { dg-line l_loop_i1 } ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i1 } ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i1 } @@ -32,11 +31,4 @@ program reduction !$acc end kernels if (red .ne. n) stop 1 - -contains - - subroutine make_addressable (v) - integer :: v ! by reference - end subroutine make_addressable - end program reduction |