From c43cb355f25dd22133d15819bd6ec03d3d3939fd Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 16 Mar 2022 14:19:41 +0100
Subject: [PATCH] Enhance further testcases to verify Openacc 'kernels'
decomposition
gcc/testsuite/
* c-c++-common/goacc-gomp/nesting-1.c: Enhance.
* c-c++-common/goacc/kernels-loop-g.c: Likewise.
* c-c++-common/goacc/nesting-1.c: Likewise.
* gcc.dg/goacc/nested-function-1.c: Likewise.
* gfortran.dg/goacc/common-block-3.f90: Likewise.
* gfortran.dg/goacc/nested-function-1.f90: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c:
Enhance.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c: Likewise.
* testsuite/libgomp.oacc-fortran/if-1.f90: Likewise.
---
.../c-c++-common/goacc-gomp/nesting-1.c | 7 ++-
.../c-c++-common/goacc/kernels-loop-g.c | 3 ++
gcc/testsuite/c-c++-common/goacc/nesting-1.c | 18 +++++--
.../gcc.dg/goacc/nested-function-1.c | 22 ++++++++
.../gfortran.dg/goacc/common-block-3.f90 | 18 +++++--
.../gfortran.dg/goacc/nested-function-1.f90 | 10 ++++
.../acc_prof-kernels-1.c | 19 +++++--
.../kernels-loop-g.c | 3 ++
.../testsuite/libgomp.oacc-fortran/if-1.f90 | 51 ++++++++++++++++++-
9 files changed, 134 insertions(+), 17 deletions(-)
@@ -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;
}
@@ -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" } */
@@ -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)
;
}
@@ -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)
@@ -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
@@ -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 }
@@ -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;
}
@@ -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"
@@ -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
--
2.34.1