diff options
9 files changed, 134 insertions, 17 deletions
diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c index 39b92712b31..51c5e359f29 100644 --- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c +++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c @@ -1,3 +1,5 @@ +/* { dg-additional-options "--param=openacc-kernels=decompose" } + /* { dg-additional-options "-fopt-info-omp-note" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } @@ -21,10 +23,11 @@ void f_acc_kernels (void) { #pragma acc kernels - /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } .-1 } - { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } .-2 } */ + /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } .-1 } */ { int i; + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} "" { target c } .+3 } + { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} "" { target c++ } .+1 } */ #pragma omp atomic write i = 0; } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-loop-g.c b/gcc/testsuite/c-c++-common/goacc/kernels-loop-g.c index 73b469d7061..5bdaa40b02c 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-loop-g.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-loop-g.c @@ -1,5 +1,8 @@ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-O2" } */ /* { dg-additional-options "-g" } */ +/*TODO PR100400 { dg-additional-options -fcompare-debug } */ /* { dg-additional-options "-fdump-tree-parloops1-all" } */ /* { dg-additional-options "-fdump-tree-optimized" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-1.c index 83cbff767a4..8c3d1adc785 100644 --- a/gcc/testsuite/c-c++-common/goacc/nesting-1.c +++ b/gcc/testsuite/c-c++-common/goacc/nesting-1.c @@ -1,3 +1,5 @@ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-all-omp" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } */ @@ -32,11 +34,13 @@ void f_acc_kernels (void) { #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i\.[0-9]+' 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_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\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 2; ++i) ; } @@ -63,15 +67,17 @@ f_acc_data (void) } #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } l_compute$c_compute } */ ; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i\.[0-9]+' 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_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\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 2; ++i) ; } @@ -102,15 +108,17 @@ f_acc_data (void) } #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } l_compute$c_compute } */ ; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */ + /* { dg-note {variable 'i\.[0-9]+' 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_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\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */ /* { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_loop_i$c_loop_i } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */ for (i = 0; i < 2; ++i) ; } diff --git a/gcc/testsuite/gcc.dg/goacc/nested-function-1.c b/gcc/testsuite/gcc.dg/goacc/nested-function-1.c index c34bcb0d601..2e48410b39d 100644 --- a/gcc/testsuite/gcc.dg/goacc/nested-function-1.c +++ b/gcc/testsuite/gcc.dg/goacc/nested-function-1.c @@ -2,6 +2,8 @@ /* See gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 for the Fortran version. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-all-omp" } */ /* { dg-additional-options "--param=openacc-privatization=noisy" } @@ -42,6 +44,11 @@ int main () #pragma acc kernels loop /* { dg-line l_compute_loop[incr c_compute_loop] } */ \ gang(num:local_arg) worker(local_arg) vector(local_arg) \ wait async(local_arg) + /* { dg-note {OpenACC 'kernels' decomposition: variable 'local_arg' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } + { dg-note {variable 'local_arg' made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'local_arg\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'local_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-note {variable 'local_i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute_loop$c_compute_loop } */ for (local_i = 0; local_i < N; ++local_i) @@ -61,6 +68,11 @@ int main () #pragma acc kernels loop /* { dg-line l_compute_loop[incr c_compute_loop] } */ \ gang(static:local_arg) worker(local_arg) vector(local_arg) \ wait(local_arg, local_arg + 1, local_arg + 2) async + /* { dg-note {OpenACC 'kernels' decomposition: variable 'local_arg' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } + { dg-note {variable 'local_arg' already made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'local_arg\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'local_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-note {variable 'local_i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute_loop$c_compute_loop } */ for (local_i = 0; local_i < N; ++local_i) @@ -87,6 +99,11 @@ int main () #pragma acc kernels loop /* { dg-line l_compute_loop[incr c_compute_loop] } */ \ gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) \ wait async(nonlocal_arg) + /* { dg-note {OpenACC 'kernels' decomposition: variable 'nonlocal_arg' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } + { dg-note {variable 'nonlocal_arg' made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'nonlocal_arg\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'nonlocal_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-note {variable 'nonlocal_i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute_loop$c_compute_loop } */ for (nonlocal_i = 0; nonlocal_i < N; ++nonlocal_i) @@ -106,6 +123,11 @@ int main () #pragma acc kernels loop /* { dg-line l_compute_loop[incr c_compute_loop] } */ \ gang(static:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) \ wait(nonlocal_arg, nonlocal_arg + 1, nonlocal_arg + 2) async + /* { dg-note {OpenACC 'kernels' decomposition: variable 'nonlocal_arg' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } + { dg-note {variable 'nonlocal_arg' already made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'nonlocal_arg\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ + /* { dg-note {variable 'nonlocal_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-note {variable 'nonlocal_i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute_loop$c_compute_loop } */ for (nonlocal_i = 0; nonlocal_i < N; ++nonlocal_i) diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 index 9dbfa4cd2f0..6f08d7eb8d5 100644 --- a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 @@ -1,5 +1,7 @@ ! { dg-options "-fopenacc -fdump-tree-omplower" } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-omp-all" } ! { dg-additional-options "--param=openacc-privatization=noisy" } @@ -28,7 +30,11 @@ program main a(i) = b(i) + c end do !$acc kernels ! { dg-line l2 } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l2 } + ! { dg-note {variable 'i' made addressable} {} { target *-*-* } l2 } + ! { dg-note {variable 'c\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l2 } ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l2 } + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } do i = 1, n x(i) = y(i) + c end do @@ -39,10 +45,14 @@ end program main ! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(tofrom:b \\\[len: 400\\\]\\\)" 1 "omplower" } } ! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\)" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(tofrom:x \\\[len: 400\\\]\\)" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(tofrom:y \\\[len: 400\\\]\\\)" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:i \\\[len: 4\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(tofrom:x \\\[len: 400\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:x \\\[len: 400\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(tofrom:y \\\[len: 400\\\]\\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\]\\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:c \\\[len: 4\\\]\\)" 1 "omplower" } } ! Expecting no mapping of un-referenced common-blocks variables diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 index 50fd0c82e14..c631f90b27f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 @@ -1,6 +1,8 @@ ! Exercise nested function decomposition, gcc/tree-nested.c. ! See gcc/testsuite/gcc.dg/goacc/nested-function-1.c for the C version. +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-all-omp" } ! { dg-additional-options "--param=openacc-privatization=noisy" } @@ -44,6 +46,8 @@ contains !$acc kernels loop & !$acc gang(num:local_arg) worker(local_arg) vector(local_arg) & !$acc wait async(local_arg) ! { dg-line l_compute_loop[incr c_compute_loop] } + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } + ! { dg-note {variable 'local_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'local_i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'local_i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } @@ -65,6 +69,8 @@ contains !$acc kernels loop & !$acc gang(static:local_arg) worker(local_arg) vector(local_arg) & !$acc wait(local_arg, local_arg + 1, local_arg + 2) async ! { dg-line l_compute_loop[incr c_compute_loop] } + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } + ! { dg-note {variable 'local_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'local_i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'local_i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } @@ -95,6 +101,8 @@ contains !$acc kernels loop & !$acc gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) & !$acc wait async(nonlocal_arg) ! { dg-line l_compute_loop[incr c_compute_loop] } + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } + ! { dg-note {variable 'nonlocal_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'nonlocal_i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'nonlocal_i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } @@ -116,6 +124,8 @@ contains !$acc kernels loop & !$acc gang(static:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) & !$acc wait(nonlocal_arg, nonlocal_arg + 1, nonlocal_arg + 2) async ! { dg-line l_compute_loop[incr c_compute_loop] } + ! { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } + ! { dg-note {variable 'nonlocal_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'nonlocal_i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'nonlocal_i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c index c82a7edbfa0..2c853971474 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c @@ -1,5 +1,7 @@ /* Test dispatch of events to callbacks. */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-fopt-info-omp-all" } { dg-additional-options "-foffload=-fopt-info-omp-all" } */ @@ -74,7 +76,7 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e assert (prof_info->device_type == acc_device_type); assert (prof_info->device_number == acc_device_num); assert (prof_info->thread_id == -1); - assert (prof_info->async == acc_async_sync); + assert (prof_info->async == acc_async_noval); assert (prof_info->async_queue == prof_info->async); assert (prof_info->src_file == NULL); assert (prof_info->func_name == NULL); @@ -181,10 +183,13 @@ int main() #define N 100 int x[N]; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not 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 } */ { + /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */ for (int i = 0; i < N; ++i) x[i] = i * i; } @@ -208,11 +213,14 @@ int main() int x[N]; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ \ num_gangs (30) num_workers (3) vector_length (5) - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not 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-warning {using 'vector_length \(32\)', ignoring 5} {} { target { __OPTIMIZE__ && openacc_nvidia_accel_selected } } 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 } */ for (int i = 0; i < N; ++i) x[i] = i * i; } @@ -236,11 +244,14 @@ int main() int x[N]; #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */ \ num_gangs (num_gangs) num_workers (num_workers) vector_length (vector_length) - /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not 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-warning {using 'vector_length \(32\)', ignoring runtime setting} {} { target { __OPTIMIZE__ && openacc_nvidia_accel_selected } } 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 } */ for (int i = 0; i < N; ++i) x[i] = i * i; } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c index 88258be16bd..e513946ea71 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c @@ -1,3 +1,6 @@ +/* { dg-additional-options "--param=openacc-kernels=decompose" } */ + /* { dg-additional-options "-g" } */ +/*TODO PR100400 { dg-additional-options -fcompare-debug } */ #include "kernels-loop.c" diff --git a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 index 3c4d9a6efb7..c6d67647d4a 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 @@ -1,6 +1,8 @@ ! { dg-do run } ! { dg-additional-options "-cpp" } +! { dg-additional-options "--param=openacc-kernels=decompose" } + ! { dg-additional-options "-fopt-info-note-omp" } ! { dg-additional-options "-foffload=-fopt-info-note-omp" } @@ -490,6 +492,9 @@ program main a(:) = 4.0 !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (1 == 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -513,6 +518,9 @@ program main a(:) = 16.0 !$acc kernels if (0 == 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -530,6 +538,9 @@ program main a(:) = 8.0 !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (one == 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -553,6 +564,9 @@ program main a(:) = 22.0 !$acc kernels if (zero == 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -570,6 +584,9 @@ program main a(:) = 16.0 !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (.TRUE.) ! { 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -593,6 +610,9 @@ program main a(:) = 76.0 !$acc kernels if (.FALSE.) ! { 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -612,6 +632,9 @@ program main nn = 1 !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (nn == 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -637,6 +660,9 @@ program main nn = 0 !$acc kernels if (nn == 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -656,6 +682,9 @@ program main nn = 1 !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if ((nn + nn) > 0) ! { 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -681,6 +710,9 @@ program main nn = 0; !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if ((nn + nn) > 0) ! { 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -698,6 +730,9 @@ program main a(:) = 91.0 !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (-2 > 0) ! { 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -715,6 +750,9 @@ program main a(:) = 43.0 !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (one == 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -738,6 +776,9 @@ program main a(:) = 87.0 !$acc kernels if (one == 0) ! { 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } if (acc_on_device (acc_device_host) .eqv. .TRUE.) then @@ -818,7 +859,10 @@ program main !$acc data copyin (a(1:N)) copyout (b(1:N)) if (1 == 1) ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } - !$acc kernels present (a(1:N)) + !$acc kernels present (a(1:N)) ! { 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N b(i) = a(i) end do @@ -862,7 +906,10 @@ program main !$acc data copyout (b(1:N)) if (1 == 1) ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } - !$acc kernels present (a(1:N)) present (b(1:N)) + !$acc kernels present (a(1:N)) present (b(1:N)) ! { 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-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } do i = 1, N b(i) = a(i) end do |