[PATCH][openacc] Disable pass_thread_jumps for IFN_UNIQUE

[Tom de Vries]

Hi,

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.

OK for trunk?

Thanks,
- Tom

[openacc] Disable pass_thread_jumps for IFN_UNIQUE

2019-06-12  Tom de Vries  <tdevr...@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.

---
 gcc/tree-ssa-threadbackward.c                      |  5 ++++
 .../testsuite/libgomp.oacc-c-c++-common/pr90009.c  | 34 ++++++++++++++++++++++
 2 files changed, 39 insertions(+)

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/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;
+}