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

[hsa] Avoid segfault in hsa switch expansion

Browse source 
2016-05-23  Martin Jambor  <mjambor@suse.cz>

	* hsa-gen.c (gen_hsa_insns_for_switch_stmt): Create an empty
	default block if a PHI node in the original one would be resized.

libgomp/
	* testsuite/libgomp.hsa.c/switch-sbr-2.c: New test.

From-SVN: r236585
This commit is contained in:
Martin Jambor 2016-05-23 13:45:13 +02:00 • committed by Martin Jambor
parent e569db5fb5
commit a50575432b
4 changed files with 74 additions and 0 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

5
gcc/ChangeLog
View file

@ -1,3 +1,8 @@
2016-05-23 Martin Jambor <mjambor@suse.cz>
* hsa-gen.c (gen_hsa_insns_for_switch_stmt): Create an empty
default block if a PHI node in the original one would be resized.
2016-05-23 Venkataramanan Kumar <venkataramanan.kumar@amd.com>
PR tree-optimization/58135

6
gcc/hsa-gen.c
View file

@ -3482,6 +3482,12 @@ gen_hsa_insns_for_switch_stmt (gswitch *s, hsa_bb *hbb)
basic_block default_label_bb = label_to_block_fn (func,
CASE_LABEL (default_label));
if (!gimple_seq_empty_p (phi_nodes (default_label_bb)))
{
default_label_bb = split_edge (find_edge (e->dest, default_label_bb));
hsa_init_new_bb (default_label_bb);
}
make_edge (e->src, default_label_bb, EDGE_FALSE_VALUE);
hsa_cfun->m_modified_cfg = true;

4
libgomp/ChangeLog
View file

@ -1,3 +1,7 @@
2016-05-23 Martin Jambor <mjambor@suse.cz>
* testsuite/libgomp.hsa.c/switch-sbr-2.c: New test.
2016-05-17 Chung-Lin Tang <cltang@codesourcery.com>
* oacc-init.c (acc_init): Remove !cached_base_dev condition on call

59
libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c Normal file
View file

@ -0,0 +1,59 @@
/* { dg-additional-options "-fno-tree-switch-conversion" } */
#pragma omp declare target
int
foo (unsigned a)
{
switch (a)
{
case 1 ... 5:
return 1;
case 9 ... 11:
return a + 3;
case 12 ... 13:
return a + 3;
default:
return 44;
}
}
#pragma omp end declare target
#define s 100
void __attribute__((noinline, noclone))
verify(int *a)
{
if (a[0] != 44)
__builtin_abort ();
for (int i = 1; i <= 5; i++)
if (a[i] != 1)
__builtin_abort ();
for (int i = 6; i <= 8; i++)
if (a[i] != 44)
__builtin_abort ();
for (int i = 9; i <= 13; i++)
if (a[i] != i + 3)
__builtin_abort ();
for (int i = 14; i < s; i++)
if (a[i] != 44)
__builtin_abort ();
}
int main(int argc)
{
int array[s];
#pragma omp target
{
for (int i = 0; i < s; i++)
{
int v = foo (i);
array[i] = v;
}
}
verify (array);
return 0;
}