From: Tom de Vries <tdevries@suse.de>
Date: Sat, 15 Jun 2019 07:06:19 +0000 (+0000)
Subject: [openacc] Disable pass_thread_jumps for IFN_UNIQUE
X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=120a01d160cd47b4276507dc8f6c1ab571a9c006;p=gcc.git

[openacc] Disable pass_thread_jumps for IFN_UNIQUE

If we compile the openacc testcase with -fopenacc -O2, we run into a SIGSEGV
or assert.  The root cause for this is that pass_thread_jumps breaks the
invariant that OACC_FORK and OACC_JOIN mark the start and end of a
single-entry-single-exit region.

Fix this by bailing out when encountering an IFN_UNIQUE in
thread_jumps::profitable_jump_thread_path.

Bootstrapped and reg-tested on x86_64.
Build and reg-tested libgomp on x86_64 with nvptx accelerator.

2019-06-15  Tom de Vries  <tdevries@suse.de>

	PR tree-optimization/90009
	* tree-ssa-threadbackward.c (thread_jumps::profitable_jump_thread_path):
	Return NULL if bb contains IFN_UNIQUE.

	* testsuite/libgomp.oacc-c-c++-common/pr90009.c: New test.

From-SVN: r272321
---

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 21b8e905eeb..4f75fb12e26 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,9 @@
+2019-06-15  Tom de Vries  <tdevries@suse.de>
+
+	PR tree-optimization/90009
+	* tree-ssa-threadbackward.c (thread_jumps::profitable_jump_thread_path):
+	Return NULL if bb contains IFN_UNIQUE.
+
 2019-06-14  Segher Boessenkool  <segher@kernel.crashing.org>
 
 	* config/rs6000/rs6000.md (CCEITHER): New define_mode_iterator.
diff --git a/gcc/tree-ssa-threadbackward.c b/gcc/tree-ssa-threadbackward.c
index 81dc05dc831..1ff870ad00b 100644
--- a/gcc/tree-ssa-threadbackward.c
+++ b/gcc/tree-ssa-threadbackward.c
@@ -261,6 +261,11 @@ thread_jumps::profitable_jump_thread_path (basic_block bbi, tree name,
 	       gsi_next_nondebug (&gsi))
 	    {
 	      gimple *stmt = gsi_stmt (gsi);
+	      if (gimple_call_internal_p (stmt, IFN_UNIQUE))
+		{
+		  m_path.pop ();
+		  return NULL;
+		}
 	      /* Do not count empty statements and labels.  */
 	      if (gimple_code (stmt) != GIMPLE_NOP
 		  && !(gimple_code (stmt) == GIMPLE_ASSIGN
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 35fe54851f9..9a1fcff1fa3 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,8 @@
+2019-06-15  Tom de Vries  <tdevries@suse.de>
+
+	PR tree-optimization/90009
+	* testsuite/libgomp.oacc-c-c++-common/pr90009.c: New test.
+
 2019-06-13  Feng Xue  <fxue@os.amperecomputing.com>
 
 	PR tree-optimization/89713
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr90009.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr90009.c
new file mode 100644
index 00000000000..58d1039dd8d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr90009.c
@@ -0,0 +1,34 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+
+#define N 100
+
+int data[N];
+
+int
+main (void)
+{
+  int n = N, b = 3;
+#pragma acc parallel num_workers(2)
+  {
+    int c;
+    if (n)
+      c = 0;
+    else
+      c = b;
+
+#pragma acc loop worker
+    for (int i = 0; i < n; i++)
+      data[i] = 1;
+
+    if (c)
+      data[0] = 2;
+  }
+
+  for (int i = 0; i < n; i++)
+    if (data[i] != 1)
+      abort ();
+
+  return 0;
+}