Hi!
On 2022-03-12T15:54:31+0100, I wrote:
> On 2022-03-01T17:46:20+0100, I wrote:
>> On 2022-01-13T10:54:16+0100, I wrote:
>>> On 2019-05-08T14:51:57+0100, Julian Brown <julian@codesourcery.com> wrote:
>>>> - The "addressable" bit is set during the kernels conversion pass for
>>>> variables that have "create" (alloc) clauses created for them in the
>>>> synthesised outer data region (instead of in the front-end, etc.,
>>>> where it can't be done accurately). Such variables actually have
>>>> their address taken during transformations made in a later pass
>>>> (omp-low, I think), but there's a phase-ordering problem that means
>>>> the flag should be set earlier.
>>>
>>> The actual issue is a bit different, but yes, there is a problem.
>>> The related ICE has also been reported as <https://gcc.gnu.org/PR100280>
>>> "ICE in lower_omp_target, at omp-low.c:12287". [...]
>
> We've resolved all such known ICEs -- but still have open
> <https://gcc.gnu.org/PR104892> "OpenACC 'kernels' decomposition:
> wrong-code cases unless manually making certain variables addressable".
> This is avoided by:
>
>> workaround patches like
>> we have on the og11 development branch:
>> - "Avoid introducing 'create' mapping clauses for loop index variables in kernels regions",
>> - "Run all kernels regions with GOMP_MAP_FORCE_TOFROM mappings synchronously",
>> - "Fix for is_gimple_reg vars to 'data kernels'"
>
> ..., but the misbehavior is visible without the workaround patches, for
> example on the master branch.
..., and to avoid the need for those (thus, please remove for
the upcoming og12 branch, Kwok), pushed to master branch
commit a07b8f4fb756484893b5612cbe9410970dc76db9
"OpenACC 'kernels' decomposition: resolve wrong-code cases unless
manually making certain variables addressable [PR100280, PR104892]",
see attached.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
From a07b8f4fb756484893b5612cbe9410970dc76db9 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 11 Mar 2022 22:31:51 +0100
Subject: [PATCH] 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.
---
gcc/omp-oacc-kernels-decompose.cc | 1 -
.../c-c++-common/goacc/kernels-decompose-2.c | 32 +++++++++++++----
.../goacc/kernels-decompose-pr100400-1-1.c | 2 ++
.../goacc/kernels-decompose-pr100400-1-2.c | 2 ++
.../goacc/kernels-decompose-pr100400-1-3.c | 2 ++
.../goacc/kernels-decompose-pr100400-1-4.c | 2 ++
.../goacc/kernels-decompose-pr104061-1-1.c | 2 ++
.../goacc/kernels-decompose-pr104061-1-2.c | 2 ++
.../goacc/kernels-decompose-pr104061-1-3.c | 2 ++
.../goacc/kernels-decompose-pr104061-1-4.c | 2 ++
.../goacc/kernels-decompose-pr104132-1.c | 2 ++
.../goacc/kernels-decompose-pr104133-1.c | 2 ++
.../goacc/kernels-decompose-pr104774-1.c | 2 ++
.../gfortran.dg/goacc/classify-kernels.f95 | 2 ++
.../gfortran.dg/goacc/kernels-decompose-2.f95 | 26 ++++++++++++--
.../libgomp.oacc-c-c++-common/declare-vla.c | 4 +++
.../libgomp.oacc-c-c++-common/default-1.c | 5 ++-
.../kernels-decompose-1.c | 16 ++++-----
.../kernels-reduction-1.c | 5 ++-
.../libgomp.oacc-c-c++-common/parallel-dims.c | 34 +++++++++++++------
.../libgomp.oacc-fortran/asyncwait-1.f90 | 6 ++++
.../kernels-reduction-1.f90 | 12 ++-----
22 files changed, 120 insertions(+), 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
--
2.34.1