From a50575432b41b3bc3b0a14cb2e6e53881941a99f Mon Sep 17 00:00:00 2001 From: Martin Jambor Date: Mon, 23 May 2016 13:45:13 +0200 Subject: [PATCH] [hsa] Avoid segfault in hsa switch expansion 2016-05-23 Martin Jambor * 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 --- gcc/ChangeLog | 5 ++ gcc/hsa-gen.c | 6 ++ libgomp/ChangeLog | 4 ++ .../testsuite/libgomp.hsa.c/switch-sbr-2.c | 59 +++++++++++++++++++ 4 files changed, 74 insertions(+) create mode 100644 libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 56ccb0ef7c9..4218dab37bb 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2016-05-23 Martin Jambor + + * 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 PR tree-optimization/58135 diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c index 697d5997519..cf7d434bce4 100644 --- a/gcc/hsa-gen.c +++ b/gcc/hsa-gen.c @@ -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; diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 3f2fe5a4df9..6a390879ff1 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,7 @@ +2016-05-23 Martin Jambor + + * testsuite/libgomp.hsa.c/switch-sbr-2.c: New test. + 2016-05-17 Chung-Lin Tang * oacc-init.c (acc_init): Remove !cached_base_dev condition on call diff --git a/libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c b/libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c new file mode 100644 index 00000000000..06990d1c2c0 --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c @@ -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; +} -- 2.30.2