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 <rguenther@suse.de>
From-SVN: r259346

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index c36052dce55c0b4d1e988a1577c41d4d3a02ac44..798469f28a4004483eb4255b8e6f727306a52791 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,12 @@
+2018-04-12  Cesar Philippidis  <cesar@codesourcery.com>
+           Richard Biener  <rguenther@suse.de>
+
+       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  <rguenther@suse.de>
 
        PR lto/85371

diff --git a/gcc/lto-streamer-out.c b/gcc/lto-streamer-out.c
index 1d2ab9757f1a3352f03f03028e59358a4f210683..fd6788a69b050c8357a9a8ab83319c659d590b24 100644 (file)
--- 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 bb204906ea64efa94a78b42bbc606c0e61f29192..c7d30ea39641e24bc618de6c186862aaf74c2841 100644 (file)
--- 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 03099159db0d1c923ac2d85555b96bcf1abfa911..9011b656a523e2e5ed4065f1caa0e5829f94e238 100644 (file)
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,9 @@
+2018-04-12  Cesar Philippidis  <cesar@codesourcery.com>
+
+       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  <jakub@redhat.com>
 
        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 (file)
index 0000000..5910b57
--- /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 (file)
index 0000000..878d8a8
--- /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