[hsa] Avoid segfault in hsa switch expansion

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.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@236585 138bc75d-0d04-0410-961f-82ee72b054a4

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 56ccb0ef7c9e27a6543732c0cd325799be1a3eac..4218dab37bb9a77ef0313518ccb468576b3ad411 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -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

diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c
index 697d5997519e17e5dd270e81701c11324e9865ca..cf7d434bce406f554ce45d35033362f47765f9f5 100644 (file)
--- 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 3f2fe5a4df9b6629ce27837532753cbaabadefe4..6a390879ff160713f6057019d9355ed2b9326fc8 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -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

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 (file)
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;
+}