Hi,
when we expand a switch statement to a SBR HSAIL instruction, we must
be careful when changing the CFG and avoid adding a new predecessor to
the default-value basic block if it has PHI nodes because we do not
provide values in the PHIs for the new edge. We can avoid it by
splitting the edge to the default-value block and having the new edge
lead to this new block, which is what the patch below does.
Bootstrapped and tested on x86_64-linux on trunk and the gcc-6 branch.
I'll commit it to both momentarily.
Thanks,
Martin
2016-05-20 Martin Jambor <mjam...@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.
---
gcc/hsa-gen.c | 6 +++
libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c | 59 ++++++++++++++++++++++++++
2 files changed, 65 insertions(+)
create mode 100644 libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c
diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c
index 697d599..cf7d434 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/testsuite/libgomp.hsa.c/switch-sbr-2.c
b/libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c
new file mode 100644
index 0000000..06990d1
--- /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.8.2
