[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

2 files changed, 39 insertions, 0 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 35fe548..9a1fcff1f 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 0000000..58d1039
--- /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;
+}