https://gcc.gnu.org/g:6505ad1b940441074b7556048907941a26cea495
commit 6505ad1b940441074b7556048907941a26cea495 Author: Tobias Burnus <[email protected]> Date: Fri Oct 21 15:31:25 2022 +0200 omp-oacc-kernels-decompose.cc: fix -fcompare-debug with GIMPLE_DEBUG GIMPLE_DEBUG were put in a parallel region of its own, which is not only pointless but also breaks -fcompare-debug. With this commit, they are handled like simple assignments: those placed are places into the same body as the loop such that only one parallel region remains as without debugging. This fixes the existing testcase libgomp.oacc-c-c++-common/kernels-loop-g.c. Note: GIMPLE_DEBUG are only accepted with -fcompare-debug; if they appear otherwise, decompose_kernels_region_body rejects them with a sorry (unchanged). Also note that there are still many xfailed tests in the c-c++-common/goacc/kernels-decompose-pr* testcases that were added in mainline commit c14ea6a72fb1ae66e3d32ac8329558497c6e4403. gcc/ChangeLog * omp-oacc-kernels-decompose.cc (top_level_omp_for_in_stmt, decompose_kernels_region_body): Handle GIMPLE_DEBUG like simple assignment. gcc/testsuite/ChangeLog * c-c++-common/goacc/kernels-decompose-pr103836-1-2.c: Adjust xfails. * c-c++-common/goacc/kernels-decompose-pr103836-1-3.c: Likewise. * c-c++-common/goacc/kernels-decompose-pr103836-1-4.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. Co-Authored-By: Sandra Loosemore <[email protected]> Diff: --- gcc/omp-oacc-kernels-decompose.cc | 5 +++-- .../goacc/kernels-decompose-pr103836-1-2.c | 2 +- .../goacc/kernels-decompose-pr103836-1-3.c | 4 ++-- .../goacc/kernels-decompose-pr103836-1-4.c | 4 ++-- .../goacc/kernels-decompose-pr104061-1-2.c | 18 +++++++++--------- .../goacc/kernels-decompose-pr104061-1-3.c | 12 ++++++------ .../goacc/kernels-decompose-pr104061-1-4.c | 12 ++++++------ 7 files changed, 29 insertions(+), 28 deletions(-) diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc index fc3a3b352d3a..5d7289ade75d 100644 --- a/gcc/omp-oacc-kernels-decompose.cc +++ b/gcc/omp-oacc-kernels-decompose.cc @@ -120,7 +120,8 @@ top_level_omp_for_in_stmt (gimple *stmt) for (gsi = gsi_start (body); !gsi_end_p (gsi); gsi_next (&gsi)) { gimple *body_stmt = gsi_stmt (gsi); - if (gimple_code (body_stmt) == GIMPLE_ASSIGN) + if (gimple_code (body_stmt) == GIMPLE_ASSIGN + || gimple_code (body_stmt) == GIMPLE_DEBUG) continue; else if (gimple_code (body_stmt) == GIMPLE_OMP_FOR && gsi_one_before_end_p (gsi)) @@ -1363,7 +1364,7 @@ decompose_kernels_region_body (gimple *kernels_region, tree kernels_clauses) = (gimple_code (stmt) == GIMPLE_ASSIGN && TREE_CODE (gimple_assign_lhs (stmt)) == VAR_DECL && DECL_ARTIFICIAL (gimple_assign_lhs (stmt))); - if (!is_simple_assignment) + if (!is_simple_assignment && gimple_code (stmt) != GIMPLE_DEBUG) only_simple_assignments = false; } } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-2.c index 83690b6d1abc..a3afb79b753f 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-2.c @@ -16,7 +16,7 @@ f_acc_kernels (void) #pragma acc kernels /* { dg-line l_compute1 } */ /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } l_compute1 } */ { - /* { 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} { target *-*-* } .-1 } */ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop /* { dg-line l_loop_i1 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-3.c index 35892a01d642..8cbf69f301c8 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-3.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-3.c @@ -1,7 +1,7 @@ /* { 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-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} "" { target *-*-* } 0 } */ /* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ /* { dg-additional-options "-fopt-info-all-omp" } */ @@ -17,7 +17,7 @@ f_acc_kernels (void) #pragma acc kernels /* { dg-line l_compute1 } */ /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ { - /* { 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} { target *-*-* } .-1 } */ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop /* { dg-line l_loop_i1 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-4.c index 549ad5dc2916..507e73cda31c 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-4.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr103836-1-4.c @@ -1,7 +1,7 @@ /* { 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-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} "" { target *-*-* } 0 } */ /* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ /* { dg-additional-options "-fopt-info-all-omp" } */ @@ -17,7 +17,7 @@ f_acc_kernels (void) #pragma acc kernels /* { dg-line l_compute1 } */ /* { dg-note {variable 'i\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute1 } */ { - /* { 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} { target *-*-* } .-1 } */ /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop /* { dg-line l_loop_i1 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c index 4d7cbb04f4b2..6865a5ca4c92 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-2.c @@ -14,23 +14,23 @@ int arr_0; void foo (void) { - /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail *-*-* } .+1 } */ + /* { dg-bogus {sorry, unimplemented: 'gimple_debug' not yet supported} TODO { xfail c++ } .+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 } */ + /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { target *-*-* } l_compute1 } */ + /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { target *-*-* } l_compute1 } */ + /* { dg-note {variable 'arr_0\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } l_compute1 } */ { - /* { 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 } */ + /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { target c++ } .-1 } + { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { target c } .+1 } */ int k; /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ #pragma acc loop /* { dg-line l_loop_k1 } */ - /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_loop_k1 } */ - /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail *-*-* } l_loop_k1 } */ - /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail *-*-* } l_loop_k1 } */ + /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++} l_loop_k1 } */ + /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { xfail c++ } l_loop_k1 } */ + /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { xfail c++ } l_loop_k1 } */ for (k = 0; k < 2; k++) arr_0 += k; } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c index 70c2ac5b5312..197cee339f02 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-3.c @@ -1,7 +1,7 @@ /* { 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 *-*-* } 0 } */ + { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} "" { target *-*-* } 0 } */ /* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ /* { dg-additional-options "-fopt-info-all-omp" } */ @@ -19,13 +19,13 @@ 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-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 } */ + /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { target *-*-* } l_compute1 } */ + /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { target *-*-* } l_compute1 } */ + /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { 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 } */ { - /* { 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 } */ + /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { target c++ } .-1 } + { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { target c } .+1 } */ int k; /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c index d1cc1a97c9f0..f82c1c72d30e 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-pr104061-1-4.c @@ -1,7 +1,7 @@ /* { 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 *-*-* } 0 } */ + { dg-bogus {error: [^\n\r]+: '-fcompare-debug' failure \(length\)} "" { target *-*-* } 0 } */ /* { dg-additional-options "-O1" } so that we may get some 'GIMPLE_DEBUG's. */ /* { dg-additional-options "-fopt-info-all-omp" } */ @@ -19,13 +19,13 @@ 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-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 } */ + /* { dg-bogus {note: OpenACC 'kernels' decomposition: variable 'k' declared in block requested to be made addressable} {w/ debug} { target *-*-* } l_compute1 } */ + /* { dg-bogus {note: variable 'k' made addressable} {w/ debug} { target *-*-* } l_compute1 } */ + /* { dg-bogus {note: variable 'k' declared in block is candidate for adjusting OpenACC privatization level} {w/ debug} { 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 } */ { - /* { 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 } */ + /* { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { target c++ } .-1 } + { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {w/ debug} { target c } .+1 } */ int k; /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
