| Message ID | 20190612175359.GA32526@delia |
|---|---|
| State | New |
| Headers | show |
| Series | [openacc] Disable pass_thread_jumps for IFN_UNIQUE | expand |
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
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; +}