Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels' decomposition [PR100400, PR103836, PR104061] (was: Decompose OpenACC 'kernels' constructs into parts, a sequence of compute constructs)
Commit Message
Hi!
On 2020-11-13T23:22:30+0100, I wrote:
> I've pushed to master branch [...] commit
> e898ce7997733c29dcab9c3c62ca102c7f9fa6eb "Decompose OpenACC 'kernels'
> constructs into parts, a sequence of compute constructs", see attached.
>
> On 2019-02-01T00:59:30+0100, I wrote:
>> There's more work to be done there, and we're aware of a number of TODO
>> items, but nevertheless: it's a good first step.
>
> That's still the case... :-)
(The pass is still disabled by default, by the way.)
We've found that 'gcc/omp-oacc-kernels-decompose.cc' is currently not at
all considerate of 'GIMPLE_DEBUG' statements -- and it's not always
straight forward how to handle these (not rocket science either; but
needs proper understanding and testing).
Actually fixing it is a separate task, but it seems prudent to at least
catch it, and document via a few test cases. OK to push
"Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels' decomposition
[PR100400, PR103836, PR104061]", 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
Comments
On Wed, Jan 19, 2022 at 11:29:18PM +0100, Thomas Schwinge wrote:
> (The pass is still disabled by default, by the way.)
>
> We've found that 'gcc/omp-oacc-kernels-decompose.cc' is currently not at
> all considerate of 'GIMPLE_DEBUG' statements -- and it's not always
> straight forward how to handle these (not rocket science either; but
> needs proper understanding and testing).
The general rule is that debug stmts shouldn't affect code generation
decisions, so when deciding what to optimize/how, they should be ignored,
and during actual transformation adjusted or worst case reset as needed.
> Actually fixing it is a separate task, but it seems prudent to at least
> catch it, and document via a few test cases. OK to push
> "Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels' decomposition
> [PR100400, PR103836, PR104061]", see attached?
> --- a/gcc/omp-oacc-kernels-decompose.cc
> +++ b/gcc/omp-oacc-kernels-decompose.cc
> @@ -1255,6 +1255,16 @@ decompose_kernels_region_body (gimple *kernels_region, tree kernels_clauses)
> gsi_next (&gsi_n);
>
> gimple *stmt = gsi_stmt (gsi);
> + if (gimple_code (stmt) == GIMPLE_DEBUG)
> + {
> + if (flag_compare_debug_opt || flag_compare_debug)
> + /* Let the usual '-fcompare-debug' analysis bail out, as
> + necessary. */
> + ;
> + else
> + sorry_at (loc, "%qs not yet supported",
> + gimple_code_name[gimple_code (stmt)]);
> + }
This is wrong. It shouldn't be dependent on flag_compare_debug* options,
those are just debugging aids to verify that -g/-g0 don't affect code
generation. With the above you'd pretend they don't, but they actually
would (with -g you'd get sorry, without it it would compile fine).
If this code is analysing whether the kernels region body should be
decomposed or not, it should be if (is_gimple_debug (stmt)) continue;
or whatever else to just ignore them (in some opts already during analysis
phase we remember they are present and something about them, but not in
a way that would actually affect the code generation decisions).
And then when actually transforming it, it depends on what transformations
are done to the variables/values referenced in the debug stmts.
gimple_debug_bind_reset_value (stmt); update_stmt (stmt); is
what resets them and can be used as last resort, it will keep saying
that it describes some var, but will say that the var is optimized out.
Jakub
Hi Jakub!
Thanks for looking into this.
On 2022-01-20T00:00:23+0100, Jakub Jelinek <jakub@redhat.com> wrote:
> On Wed, Jan 19, 2022 at 11:29:18PM +0100, Thomas Schwinge wrote:
>> (The pass is still disabled by default, by the way.)
>>
>> We've found that 'gcc/omp-oacc-kernels-decompose.cc' is currently not at
>> all considerate of 'GIMPLE_DEBUG' statements -- and it's not always
>> straight forward how to handle these (not rocket science either; but
>> needs proper understanding and testing).
>
> The general rule is that debug stmts shouldn't affect code generation
> decisions, so when deciding what to optimize/how, they should be ignored
ACK. (... and I'm confused why we didn't run into this when originally
doing the OpenACC 'kernels' decomposition work, three years ago...)
> and during actual transformation adjusted or worst case reset as needed.
That's what we need to look into, in particular: if we decompose (GIMPLE
sequence) an OpenACC 'kernels' region into parts, how to move or
otherwise handle any 'GIMPLE_DEBUG's.
>> Actually fixing it is a separate task, but it seems prudent to at least
>> catch it, and document via a few test cases. OK to push
>> "Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels' decomposition
>> [PR100400, PR103836, PR104061]", see attached?
>
>> --- a/gcc/omp-oacc-kernels-decompose.cc
>> +++ b/gcc/omp-oacc-kernels-decompose.cc
>> @@ -1255,6 +1255,16 @@ decompose_kernels_region_body (gimple *kernels_region, tree kernels_clauses)
>> gsi_next (&gsi_n);
>>
>> gimple *stmt = gsi_stmt (gsi);
>> + if (gimple_code (stmt) == GIMPLE_DEBUG)
>> + {
>> + if (flag_compare_debug_opt || flag_compare_debug)
>> + /* Let the usual '-fcompare-debug' analysis bail out, as
>> + necessary. */
>> + ;
>> + else
>> + sorry_at (loc, "%qs not yet supported",
>> + gimple_code_name[gimple_code (stmt)]);
>> + }
>
> This is wrong.
I have a different understanding what "wrong" means. ;-)
> It shouldn't be dependent on flag_compare_debug* options,
> those are just debugging aids to verify that -g/-g0 don't affect code
> generation. With the above you'd pretend they don't, but they actually
> would (with -g you'd get sorry, without it it would compile fine).
The idea there is: not all 'GIMPLE_DEBUG's are mishandled in the pass,
just some. If '-fcompare-debug' is in effect, we know that it will
detect any cases of mishandling (code generation difference), so it's
thus fine in that case to skip the coarse-grained 'sorry' here.
> If this code is analysing whether the kernels region body should be
> decomposed or not
This place here is just a convenient one, where we iterate through the
whole GIMPLE sequence.
With these things now hopfully clarified, is the attached
"Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels' decomposition
[PR100400, PR103836, PR104061]" OK to push? It's of course not the final
fix, but it at least makes obvious any current silent miscompilation, and
incremental improvement over the current status.
> it should be if (is_gimple_debug (stmt)) continue;
> or whatever else to just ignore them (in some opts already during analysis
> phase we remember they are present and something about them, but not in
> a way that would actually affect the code generation decisions).
> And then when actually transforming it, it depends on what transformations
> are done to the variables/values referenced in the debug stmts.
> gimple_debug_bind_reset_value (stmt); update_stmt (stmt); is
> what resets them and can be used as last resort, it will keep saying
> that it describes some var, but will say that the var is optimized out.
Thanks, that'll be helpful later.
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
On Thu, Jan 20, 2022 at 09:26:50AM +0100, Thomas Schwinge wrote:
> That's what we need to look into, in particular: if we decompose (GIMPLE
> sequence) an OpenACC 'kernels' region into parts, how to move or
> otherwise handle any 'GIMPLE_DEBUG's.
I admit I haven't looked at the pass except now for the toplevel comment.
It says that OpenACC constructs in the region are perhaps adjusted but
their body is unchanged, so that suggests that debug stmts inside of those
bodies should be kept as is.
Next it says that sequential code in between those loops/whatever are
put into some sequential construct, so I guess if you decide so because
of some non-debug stmts, you can just move the debug stmts into that
construct as well, including those debug stmts before the first such
non-debug stmt and debug stmts after the last such non-debug stmts.
It is not a perfect solution, because normally debug stmts before
loops would affect also what is in the loop unless overridden, but
what the pass does seems terribly destructive for debug experience anyway.
There is then another case, only debug stmts e.g. in between or before
the loops or after them and nothing else. Perhaps throwing them away at
this point is the best thing to do (but, all of this only after the pass
decides that it will change something).
Another thing is, this is apparently a very early pass, so most real
debug stmts don't exist, they are typically created later.
I'd expect you mostly see gimple_debug_begin_stmt_p stmts.
Those can be removed more easily, it doesn't mean var has this value
for the following code until stated otherwise, but it just said here was
the start of some source code statement. So, if you drop them, all that
will work worse is break some_line.
So citing from e.g. PR100400:
void foo ()
{
# DEBUG BEGIN_STMT // Outside of region, don't touch this
#pragma omp target oacc_kernels map(force_tofrom:p [len: 8])
{
int c.0;
# DEBUG BEGIN_STMT // Drop this
try
{
# DEBUG BEGIN_STMT // If p = &c; is moved somewhere, move the surrounding DEBUG BEGIN_STMTs with it
# DEBUG BEGIN_STMT
p = &c;
# DEBUG BEGIN_STMT // Up to here
#pragma acc loop independent private(c.0) private(c)
for (c.0 = 0; c.0 < 1; c.0 = c.0 + 1)
{
c = c.0;
# DEBUG BEGIN_STMT // Keep this in the body
}
}
finally
{
c = {CLOBBER};
}
}
}
If you don't have time for it right now, after deciding you are
going to transform it just gsi_remove gimple_debug_begin_stmt_p stmts
you don't know how to handle.
> With these things now hopfully clarified, is the attached
> "Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels' decomposition
> [PR100400, PR103836, PR104061]" OK to push? It's of course not the final
> fix, but it at least makes obvious any current silent miscompilation, and
> incremental improvement over the current status.
No, users really don't want to see sorry messages just because they turned
-g on their code. They might be ok with their kernels not being easily
debuggable, but they surely will not be ok with not being able to debug
the host code in the same TU.
Jakub
From 568808ef7ccc97ebeae90bc7cb1aba6bd7659b24 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 19 Jan 2022 14:04:42 +0100
Subject: [PATCH] Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels'
decomposition [PR100400, PR103836, PR104061]
Actually fixing it is a separate task, but it seems prudent to at least catch
it, and document via a few test cases.
gcc/
PR middle-end/100400
PR middle-end/103836
PR middle-end/104061
* omp-oacc-kernels-decompose.cc (decompose_kernels_region_body):
Catch 'GIMPLE_DEBUG'.
gcc/testsuite/
PR middle-end/100400
PR middle-end/103836
PR middle-end/104061
* c-c++-common/goacc/kernels-decompose-pr100400-1-1.c: New.
* c-c++-common/goacc/kernels-decompose-pr100400-1-2.c: New.
* c-c++-common/goacc/kernels-decompose-pr100400-1-3.c: New.
* c-c++-common/goacc/kernels-decompose-pr100400-1-4.c: New.
* c-c++-common/goacc/kernels-decompose-pr103836-1-1.c: New.
* c-c++-common/goacc/kernels-decompose-pr103836-1-2.c: New.
* c-c++-common/goacc/kernels-decompose-pr103836-1-3.c: New.
* c-c++-common/goacc/kernels-decompose-pr103836-1-4.c: New.
* c-c++-common/goacc/kernels-decompose-pr104061-1-1.c: New.
* c-c++-common/goacc/kernels-decompose-pr104061-1-2.c: New.
* c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: New.
* c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: New.
---
gcc/omp-oacc-kernels-decompose.cc | 10 +++++
.../goacc/kernels-decompose-pr100400-1-1.c | 33 ++++++++++++++
.../goacc/kernels-decompose-pr100400-1-2.c | 40 +++++++++++++++++
.../goacc/kernels-decompose-pr100400-1-3.c | 42 ++++++++++++++++++
.../goacc/kernels-decompose-pr100400-1-4.c | 40 +++++++++++++++++
.../goacc/kernels-decompose-pr103836-1-1.c | 26 +++++++++++
.../goacc/kernels-decompose-pr103836-1-2.c | 29 +++++++++++++
.../goacc/kernels-decompose-pr103836-1-3.c | 30 +++++++++++++
.../goacc/kernels-decompose-pr103836-1-4.c | 30 +++++++++++++
.../goacc/kernels-decompose-pr104061-1-1.c | 30 +++++++++++++
.../goacc/kernels-decompose-pr104061-1-2.c | 33 ++++++++++++++
.../goacc/kernels-decompose-pr104061-1-3.c | 43 +++++++++++++++++++
.../goacc/kernels-decompose-pr104061-1-4.c | 41 ++++++++++++++++++
13 files changed, 427 insertions(+)
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-2.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-3.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr100400-1-4.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-2.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-3.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-4.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c
@@ -1255,6 +1255,16 @@ decompose_kernels_region_body (gimple *kernels_region, tree kernels_clauses)
gsi_next (&gsi_n);
gimple *stmt = gsi_stmt (gsi);
+ if (gimple_code (stmt) == GIMPLE_DEBUG)
+ {
+ if (flag_compare_debug_opt || flag_compare_debug)
+ /* Let the usual '-fcompare-debug' analysis bail out, as
+ necessary. */
+ ;
+ else
+ sorry_at (loc, "%qs not yet supported",
+ gimple_code_name[gimple_code (stmt)]);
+ }
gimple *omp_for = top_level_omp_for_in_stmt (stmt);
bool is_unconditional_oacc_for_loop = false;
if (omp_for != NULL)
new file mode 100644
@@ -0,0 +1,33 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-g0" } */
+/* { dg-additional-options "-O1" } */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+int *p;
+
+void
+foo (void)
+{
+#pragma acc kernels
+ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ {
+ int c;
+
+ /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+ p = &c;
+
+ /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+#pragma acc loop independent
+ /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } .-2 }
+ { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } .-3 } */
+ /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } .-4 } */
+ for (c = 0; c < 1; ++c)
+ ;
+ }
+}
new file mode 100644
@@ -0,0 +1,40 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-fchecking" }
+ { dg-ice TODO { c++ } }
+ { dg-prune-output "during GIMPLE pass: omp_oacc_kernels_decompose" } */
+
+/* { dg-additional-options "-g" } */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+int *p;
+
+void
+foo (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail *-*-* } .+1 } */
+#pragma acc kernels
+ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { xfail *-*-* } .-1 } */
+ /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-2 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 }
+ { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */
+ int c;
+
+ /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { xfail *-*-* } .+1 } */
+ p = &c;
+
+ /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { xfail c++ } .+1 } */
+#pragma acc loop independent
+ /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-1 } */
+ /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { xfail *-*-* } .-2 }
+ { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { xfail *-*-* } .-3 } */
+ /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { xfail *-*-* } .-4 } */
+ for (c = 0; c < 1; ++c)
+ ;
+ }
+}
new file mode 100644
@@ -0,0 +1,42 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-fchecking" }
+ { dg-ice TODO { c++ } }
+ { dg-prune-output "during GIMPLE pass: omp_oacc_kernels_decompose" } */
+
+/* { dg-additional-options "-fcompare-debug" } -- w/o debug compiled first.
+ { dg-bogus {error: during '-fcompare-debug' recompilation} TODO { xfail c++ } 0 }
+ { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail c++ } 0 } */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+int *p;
+
+void
+foo (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */
+#pragma acc kernels
+ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 }
+ { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */
+ int c;
+
+ /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+ p = &c;
+
+ /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
+#pragma acc loop independent
+ /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } .-2 }
+ { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } .-3 } */
+ /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } .-4 } */
+ for (c = 0; c < 1; ++c)
+ ;
+ }
+}
new file mode 100644
@@ -0,0 +1,40 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-fchecking" }
+ { dg-ice TODO { c++ } }
+ { dg-prune-output "during GIMPLE pass: omp_oacc_kernels_decompose" } */
+
+/* { dg-additional-options "-g -fcompare-debug" } -- w/ debug compiled first. */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+int *p;
+
+void
+foo (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */
+#pragma acc kernels
+ /* { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { xfail c++ } .-1 } */
+ /* { dg-note {variable 'c\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-2 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 }
+ { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */
+ int c;
+
+ /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { xfail c++ } .+1 } */
+ p = &c;
+
+ /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { xfail c++ } .+1 } */
+#pragma acc loop independent
+ /* { dg-note {variable 'c\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-1 } */
+ /* { dg-note {variable 'c' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { xfail c++ } .-2 }
+ { dg-note {variable 'c' ought to be adjusted for OpenACC privatization level: 'vector'} {} { xfail c++ } .-3 } */
+ /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { xfail c++ } .-4 } */
+ for (c = 0; c < 1; ++c)
+ ;
+ }
+}
new file mode 100644
@@ -0,0 +1,26 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-g0" } */
+/* { dg-additional-options "-O1" } */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+extern int i;
+
+void
+f_acc_kernels (void)
+{
+#pragma acc kernels
+ /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ {
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-3 } */
+ for (i = 0; i < 2; ++i)
+ ;
+ }
+}
new file mode 100644
@@ -0,0 +1,29 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-g" } */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+extern int i;
+
+void
+f_acc_kernels (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail c++ } .+1 } */
+#pragma acc kernels
+ /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-1 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } */
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-1 } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } .-2 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail c++ } .-3 } */
+ for (i = 0; i < 2; ++i)
+ ;
+ }
+}
new file mode 100644
@@ -0,0 +1,30 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-fcompare-debug" } -- w/o debug compiled first.
+ { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail c++ } 0 } */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+extern int i;
+
+void
+f_acc_kernels (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } */
+#pragma acc kernels
+ /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } */
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-3 } */
+ for (i = 0; i < 2; ++i)
+ ;
+ }
+}
new file mode 100644
@@ -0,0 +1,30 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-g -fcompare-debug" } -- w/ debug compiled first.
+ { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail c++ } 0 } */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" } */
+
+extern int i;
+
+void
+f_acc_kernels (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } */
+#pragma acc kernels
+ /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 } */
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'i\.0' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-3 } */
+ for (i = 0; i < 2; ++i)
+ ;
+ }
+}
new file mode 100644
@@ -0,0 +1,30 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-g0" } */
+/* { dg-additional-options "-O1" } */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" }
+ Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
+ { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
+
+int arr_0;
+
+void
+foo (void)
+{
+#pragma acc kernels
+ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ {
+ int k;
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-3 } */
+ for (k = 0; k < 2; k++)
+ arr_0 += k;
+ }
+}
new file mode 100644
@@ -0,0 +1,33 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-g" } */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" }
+ Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
+ { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
+
+int arr_0;
+
+void
+foo (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail *-*-* } .+1 } */
+#pragma acc kernels
+ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-1 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 }
+ { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */
+ int k;
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-1 } */
+ /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-2 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail *-*-* } .-3 } */
+ for (k = 0; k < 2; k++)
+ arr_0 += k;
+ }
+}
new file mode 100644
@@ -0,0 +1,43 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-fchecking" }
+ { dg-ice TODO }
+ { dg-prune-output {D\.[0-9]+ = arr_0\.0 \+ k;} }
+ { dg-prune-output {during GIMPLE pass: lower} } */
+
+/* { dg-additional-options "-fcompare-debug" } -- w/o debug compiled first.
+ { dg-bogus {error: during '-fcompare-debug' recompilation} TODO { xfail *-*-* } 0 }
+ { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} TODO { xfail *-*-* } 0 } */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" }
+ Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
+ { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
+
+int arr_0;
+
+void
+foo (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */
+#pragma acc kernels
+ /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } .-1 } */
+ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 }
+ { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */
+ int k;
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-1 } */
+ /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ /* { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } .-3 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } .-4 } */
+ for (k = 0; k < 2; k++)
+ arr_0 += k;
+ /* { dg-bogus {error: invalid operands in binary operation} {w/ debug} { xfail *-*-* } .-1 } */
+ }
+}
new file mode 100644
@@ -0,0 +1,41 @@
+/* { dg-additional-options "--param openacc-kernels=decompose" } */
+
+/* { dg-additional-options "-fchecking" }
+ { dg-ice TODO }
+ { dg-prune-output {D\.[0-9]+ = arr_0\.0 \+ k;} }
+ { dg-prune-output {during GIMPLE pass: lower} } */
+
+/* { dg-additional-options "-g -fcompare-debug" } -- w/ debug compiled first. */
+/* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */
+
+/* { dg-additional-options "-fopt-info-all-omp" } */
+
+/* { dg-additional-options "--param=openacc-privatization=noisy" }
+ Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
+ { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
+
+int arr_0;
+
+void
+foo (void)
+{
+ /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} {} { target *-*-* } .+1 } suppressed via '-fcompare-debug'. */
+#pragma acc kernels
+ /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } .-1 } */
+ /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } .-2 } */
+ {
+ /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c++ } .-1 }
+ { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { xfail c } .+1 } */
+ int k;
+
+ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
+#pragma acc loop
+ /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-1 } */
+ /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } .-2 } */
+ /* { dg-bogus {note: variable 'k' in 'private' clause is candidate for adjusting OpenACC privatization level} {w/ debug} { xfail *-*-* } .-3 } */
+ /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail *-*-* } .-4 } */
+ for (k = 0; k < 2; k++)
+ arr_0 += k;
+ /* { dg-bogus {error: invalid operands in binary operation} {w/ debug} { xfail *-*-* } .-1 } */
+ }
+}
--
2.25.1