FreeChainXenon/gcc
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

[openacc] Disable pass_thread_jumps for IFN_UNIQUE

Browse source 
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
This commit is contained in:
Tom de Vries 2019-06-15 07:06:19 +00:00 • committed by Tom de Vries
parent 2789efe3ee
commit 120a01d160
4 changed files with 50 additions and 0 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

6
gcc/ChangeLog
View file

@ -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.

5
gcc/tree-ssa-threadbackward.c
View file

@ -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

5
libgomp/ChangeLog
View file

@ -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

34
libgomp/testsuite/libgomp.oacc-c-c++-common/pr90009.c Normal file
View file

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