[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 21b8e90..4f75fb1 100644 (file)
--- 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 81dc05d..1ff870a 100644 (file)
--- 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 35fe548..9a1fcff 100644 (file)
--- 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 (file)
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;
+}