From 120a01d160cd47b4276507dc8f6c1ab571a9c006 Mon Sep 17 00:00:00 2001 From: Tom de Vries Date: Sat, 15 Jun 2019 07:06:19 +0000 Subject: [PATCH] [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 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 --- gcc/ChangeLog | 6 ++++ gcc/tree-ssa-threadbackward.c | 5 ++++ libgomp/ChangeLog | 5 ++++ .../testsuite/libgomp.oacc-c-c++-common/pr90009.c | 34 ++++++++++++++++++++++ 4 files changed, 50 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr90009.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 21b8e90..4f75fb1 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,9 @@ +2019-06-15 Tom de Vries + + 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 * 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 --- 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..9a1fcff1f 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,8 @@ +2019-06-15 Tom de Vries + + PR tree-optimization/90009 + * testsuite/libgomp.oacc-c-c++-common/pr90009.c: New test. + 2019-06-13 Feng Xue 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 + +#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; +} -- 2.7.4