public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
From: Tom de Vries <Tom_deVries@mentor.com>
To: Jakub Jelinek <jakub@redhat.com>,
	Cesar Philippidis	<cesar@codesourcery.com>
Cc: "H.J. Lu" <hjl.tools@gmail.com>,
	Richard Biener <rguenther@suse.de>,
	"gcc-patches@gcc.gnu.org" <gcc-patches@gcc.gnu.org>
Subject: Re: [PATCH] Handle empty infinite loops in OpenACC for PR84955
Date: Mon, 16 Apr 2018 18:13:00 -0000	[thread overview]
Message-ID: <b574d8e8-586c-43fb-1e6a-5ece2af20237@mentor.com> (raw)
In-Reply-To: <20180412185840.GR8577@tucnak>

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

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).

I've filed the lto ICE (described as "the second problem" in this 
thread) as PR85422 - "[openacc] ICE at cfgloop.c:468: segfault in 
flow_loops_find".

[ When changing dg-do compile to link in the test-cases in this patch, 
we run into PR85422, which shows that in the reverted patch the 
problematic fix was actually not exercised by the test-cases. ]

Thanks,
- Tom

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

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

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 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);
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

  reply	other threads:[~2018-04-16 18:13 UTC|newest]

Thread overview: 12+ messages / expand[flat|nested]  mbox.gz  Atom feed  top
2018-04-06 13:49 Cesar Philippidis
2018-04-06 15:35 ` Jakub Jelinek
2018-04-09 11:32   ` Richard Biener
2018-04-11 19:31     ` Cesar Philippidis
2018-04-12  8:02       ` Richard Biener
2018-04-12 18:27       ` H.J. Lu
2018-04-12 18:39         ` Cesar Philippidis
2018-04-12 20:35           ` Jakub Jelinek
2018-04-16 18:13             ` Tom de Vries [this message]
2018-04-23 12:07               ` [PATCH, lto, PR85422] Fixup loops before lto write-out Tom de Vries
2018-04-23 12:11                 ` Richard Biener
2018-05-01 13:31               ` [PATCH] Handle empty infinite loops in OpenACC for PR84955 Tom de Vries

Reply instructions:

You may reply publicly to this message via plain-text email
using any one of the following methods:

* Save the following mbox file, import it into your mail client,
  and reply-to-all from there: mbox

  Avoid top-posting and favor interleaved quoting:
  https://en.wikipedia.org/wiki/Posting_style#Interleaved_style

* Reply using the --to, --cc, and --in-reply-to
  switches of git-send-email(1):

  git send-email \
    --in-reply-to=b574d8e8-586c-43fb-1e6a-5ece2af20237@mentor.com \
    --to=tom_devries@mentor.com \
    --cc=cesar@codesourcery.com \
    --cc=gcc-patches@gcc.gnu.org \
    --cc=hjl.tools@gmail.com \
    --cc=jakub@redhat.com \
    --cc=rguenther@suse.de \
    /path/to/YOUR_REPLY

  https://kernel.org/pub/software/scm/git/docs/git-send-email.html

* If your mail client supports setting the In-Reply-To header
  via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).