From 2e5efa6760314aac101b63d1a7742cff5f5e650f Mon Sep 17 00:00:00 2001 From: Cesar Philippidis Date: Thu, 12 Apr 2018 06:15:45 -0700 Subject: [PATCH] re PR middle-end/84955 (Incorrect OpenACC tile expansion) PR middle-end/84955 gcc/ * lto-streamer-out.c (output_function): Fix CFG loop state before streaming out. * omp-expand.c (expand_oacc_for): Handle calls to internal functions like regular functions. libgomp/ * testsuite/libgomp.oacc-c-c++-common/pr84955.c: New test. * testsuite/libgomp.oacc-fortran/pr84955.f90: New test. Co-Authored-By: Richard Biener From-SVN: r259346 --- gcc/ChangeLog | 9 +++++++++ gcc/lto-streamer-out.c | 8 ++++++-- gcc/omp-expand.c | 8 ++++++++ gcc/testsuite/ChangeLog | 6 ++++++ .../libgomp.oacc-c-c++-common/pr84955.c | 20 +++++++++++++++++++ .../libgomp.oacc-fortran/pr84955.f90 | 20 +++++++++++++++++++ 6 files changed, 69 insertions(+), 2 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr84955.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/pr84955.f90 diff --git a/gcc/ChangeLog b/gcc/ChangeLog index c36052dce55..798469f28a4 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,12 @@ +2018-04-12 Cesar Philippidis + Richard Biener + + PR middle-end/84955 + * lto-streamer-out.c (output_function): Fix CFG loop state before + streaming out. + * omp-expand.c (expand_oacc_for): Handle calls to internal + functions like regular functions. + 2018-04-12 Richard Biener PR lto/85371 diff --git a/gcc/lto-streamer-out.c b/gcc/lto-streamer-out.c index 1d2ab9757f1..fd6788a69b0 100644 --- a/gcc/lto-streamer-out.c +++ b/gcc/lto-streamer-out.c @@ -2084,6 +2084,9 @@ output_function (struct cgraph_node *node) /* Set current_function_decl and cfun. */ push_cfun (fn); + /* Fixup loops if required to match discovery done in the reader. */ + loop_optimizer_init (AVOID_CFG_MODIFICATIONS); + /* Make string 0 be a NULL string. */ streamer_write_char_stream (ob->string_stream, 0); @@ -2176,12 +2179,13 @@ output_function (struct cgraph_node *node) streamer_write_record_start (ob, LTO_null); output_cfg (ob, fn); - - pop_cfun (); } else streamer_write_uhwi (ob, 0); + loop_optimizer_finalize (); + pop_cfun (); + /* Create a section to hold the pickled output of this function. */ produce_asm (ob, function); diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index bb204906ea6..c7d30ea3964 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -5439,6 +5439,14 @@ expand_oacc_for (struct omp_region *region, struct omp_for_data *fd) split->flags ^= EDGE_FALLTHRU | EDGE_TRUE_VALUE; + /* Add a dummy exit for the tiled block when cont_bb is missing. */ + if (cont_bb == NULL) + { + edge e = make_edge (body_bb, exit_bb, EDGE_FALSE_VALUE); + e->probability = profile_probability::even (); + split->probability = profile_probability::even (); + } + /* Initialize the user's loop vars. */ gsi = gsi_start_bb (elem_body_bb); expand_oacc_collapse_vars (fd, true, &gsi, counts, e_offset); diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 03099159db0..9011b656a52 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,9 @@ +2018-04-12 Cesar Philippidis + + PR middle-end/84955 + * testsuite/libgomp.oacc-c-c++-common/pr84955.c: New test. + * testsuite/libgomp.oacc-fortran/pr84955.f90: New test. + 2018-04-12 Jakub Jelinek PR target/85328 diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84955.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84955.c new file mode 100644 index 00000000000..5910b57b68d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84955.c @@ -0,0 +1,20 @@ +/* { dg-do compile } */ + +int +main () +{ + int i, j; + +#pragma acc parallel loop tile(2,3) + for (i = 1; i < 10; i++) + for (j = 1; j < 10; j++) + for (;;) + ; + +#pragma acc parallel loop + for (i = 1; i < 10; i++) + for (;;) + ; + + return i + j; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr84955.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr84955.f90 new file mode 100644 index 00000000000..878d8a89f41 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr84955.f90 @@ -0,0 +1,20 @@ +! { dg-do compile } + +subroutine s + integer :: i, j + !$acc parallel loop tile(2,3) + do i = 1, 10 + do j = 1, 10 + do + end do + end do + end do + !$acc end parallel loop + + !$acc parallel loop + do i = 1, 10 + do + end do + end do + !$acc end parallel loop +end subroutine s -- 2.30.2