Re: [PATCH] Handle empty infinite loops in OpenACC for PR84955

[Tom de Vries <Tom_deVries@mentor.com>]

[-- Attachment #1: Type: text/plain, Size: 1502 bytes --]

On 04/16/2018 08:13 PM, Tom de Vries wrote:
> On 04/12/2018 08:58 PM, Jakub Jelinek wrote:
>> On Thu, Apr 12, 2018 at 11:39:43AM -0700, Cesar Philippidis wrote:
>>> Strange. I didn't observe any regressions when I tested it. But, then
>>> again, I was testing against revision
>>>
>>> r259092 | jason | 2018-04-04 09:42:55 -0700 (Wed, 04 Apr 2018) | 4 lines
>>>
>>> which is over a week old. I'll revert that patch for now, and revisit
>>> this issue in stage1.
>>
>> You should have kept the omp-expand.c chunk, that is correct and 
>> shouldn't
>> cause issues.
> 
> Committed as attached (with correct changelog entry, the previously 
> committed patch had an incorrect one).
> 

Backported to gcc-7-branch as attached.

Thanks,
- Tom

> diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
> index bb20490..c7d30ea 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);


[-- Attachment #2: 0001-backport-openacc-Fix-ICE-when-compiling-tile-loop-containing-infinite-loop.patch --]
[-- Type: text/x-patch, Size: 2383 bytes --]

backport "[openacc] Fix ICE when compiling tile loop containing infinite loop"

2018-04-16  Tom de Vries  <tom@codesourcery.com>

	backport from trunk:
	2018-04-16  Cesar Philippidis  <cesar@codesourcery.com>
		    Tom de Vries  <tom@codesourcery.com>

	PR middle-end/84955
	* omp-expand.c (expand_oacc_for): Add dummy false branch for
	tiled basic blocks without omp continue statements.

	* testsuite/libgomp.oacc-c-c++-common/pr84955.c: New test.
	* testsuite/libgomp.oacc-fortran/pr84955.f90: New test.

---
 gcc/omp-expand.c                                      |  8 ++++++++
 libgomp/testsuite/libgomp.oacc-c-c++-common/pr84955.c | 15 +++++++++++++++
 libgomp/testsuite/libgomp.oacc-fortran/pr84955.f90    | 13 +++++++++++++
 3 files changed, 36 insertions(+)

diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index a1e668d..549a5db 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -5628,6 +5628,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 = PROB_EVEN;
+	      split->probability = PROB_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/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84955.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84955.c
new file mode 100644
index 0000000..e528faa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84955.c
@@ -0,0 +1,15 @@
+/* { dg-do compile }  */
+
+int
+main (void)
+{
+  int i, j;
+
+#pragma acc parallel loop tile(2,3)
+  for (i = 1; i < 10; i++)
+    for (j = 1; j < 10; j++)
+      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 0000000..dc85865
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr84955.f90
@@ -0,0 +1,13 @@
+! { 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
+end subroutine s