From: Martin Jambor <mjambor@suse.cz>
Date: Mon, 23 May 2016 11:45:13 +0000 (+0200)
Subject: [hsa] Avoid segfault in hsa switch expansion
X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=a50575432b41b3bc3b0a14cb2e6e53881941a99f;p=gcc.git

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

From-SVN: r236585
---

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