OpenACC 'kernels' decomposition: resolve wrong-code cases unless manually making certain variables addressable [PR100280, PR104892]
Currently in OpenACC 'kernels' decomposition, there is special handling of 'GOMP_MAP_FORCE_TOFROM', documented to be done to avoid "internal compiler errors in later passes". For performance reasons, the current repetitive to/from device copying for every region is not ideal, compared to using 'present' clauses, as done for almost all other 'GOMP_MAP_*'. Also, the current special handling (incomplete, evidently) is the reason for the PR104892 misbehavior. For PR100280 etc. we've resolved all such known ICEs -- removing the special handling for 'GOMP_MAP_FORCE_TOFROM' now resolves PR104892. PR middle-end/100280 PR middle-end/104892 gcc/ * omp-oacc-kernels-decompose.cc (omp_oacc_kernels_decompose_1): Remove special handling of 'GOMP_MAP_FORCE_TOFROM'. gcc/testsuite/ * c-c++-common/goacc/kernels-decompose-2.c: Adjust. * c-c++-common/goacc/kernels-decompose-pr100400-1-1.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr100400-1-2.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr100400-1-3.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr100400-1-4.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104061-1-1.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104061-1-2.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104132-1.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104133-1.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr104774-1.c: Likewise. * gfortran.dg/goacc/classify-kernels.f95: Likewise. * gfortran.dg/goacc/kernels-decompose-2.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Adjust. * testsuite/libgomp.oacc-c-c++-common/default-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-reduction-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise. * testsuite/libgomp.oacc-fortran/asyncwait-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-reduction-1.f90: Likewise.
This commit is contained in:
Thomas Schwinge
parent
535afbd959
commit
a07b8f4fb7
22 changed files with 120 additions and 45 deletions
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
|
|
|
|||
|
|
@ -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 } */
|
||||
{
|
||||
|
|
|
|||
|
|
@ -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 } */
|
||||
{
|
||||
|
|
|
|||
|
|
@ -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 } */
|
||||
{
|
||||
|
|
|
|||
|
|
@ -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 } */
|
||||
{
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
|
|
|
|||
|
|
@ -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 } */
|
||||
|
|
|
|||
|
|
@ -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 } */
|
||||
|
|
|
|||
|
|
@ -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 } */
|
||||
|
|
|
|||
|
|
@ -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 } */
|
||||
|
|
|
|||
|
|
@ -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 } */
|
||||
|
|
|
|||
|
|
@ -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 } */
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
||||
|
|
|
|||
|
|
@ -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 } */
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
|
|
|
|||
|
|
@ -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 ();
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -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 } */
|
||||
|
|
|
|||
|
|
@ -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)
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue