[openacc] Disable pass_thread_jumps for IFN_UNIQUE

Message ID 20190612175359.GA32526@delia
State New
Headers show
Series
  • [openacc] Disable pass_thread_jumps for IFN_UNIQUE
Related show

Commit Message

Tom de Vries June 12, 2019, 5:54 p.m.
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  <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.

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

Comments

Jeff Law June 14, 2019, 8:09 p.m. | #1
On 6/12/19 11:54 AM, Tom de Vries wrote:
> 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  <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.

This really isn't a profitability test, but it's awkward to put the test
in the earlier statement scan.  So OK for the trunk.

jeff

Patch

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