Fix switch conversion in offloading functions

2018-03-26  Tom de Vries  <tom@codesourcery.com>

	PR tree-optimization/85063
	* omp-general.c (offloading_function_p): New function.  Factor out
	of ...
	* omp-offload.c (pass_omp_target_link::gate): ... here.
	* omp-general.h (offloading_function_p): Declare.
	* tree-switch-conversion.c (build_one_array): Mark CSWTCH.x variable
	with attribute omp declare target for offloading functions.

	* testsuite/libgomp.c/switch-conversion-2.c: New test.
	* testsuite/libgomp.c/switch-conversion.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/switch-conversion-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/switch-conversion.c: New test.

From-SVN: r258852

5 files changed, 141 insertions, 0 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 008b93535fb..a5a5e0631b2 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,11 @@
+2018-03-26  Tom de Vries  <tom@codesourcery.com>
+
+	PR tree-optimization/85063
+	* testsuite/libgomp.c/switch-conversion-2.c: New test.
+	* testsuite/libgomp.c/switch-conversion.c: New test.
+	* testsuite/libgomp.oacc-c-c++-common/switch-conversion-2.c: New test.
+	* testsuite/libgomp.oacc-c-c++-common/switch-conversion.c: New test.
+
 2018-03-25  Thomas Koenig  <tkoenig@gcc.gnu.org>
 
 	PR fortran/84381
diff --git a/libgomp/testsuite/libgomp.c/switch-conversion-2.c b/libgomp/testsuite/libgomp.c/switch-conversion-2.c
new file mode 100644
index 00000000000..97601dc9cbb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/switch-conversion-2.c
@@ -0,0 +1,31 @@
+/* PR tree-optimization/85063 */
+/* { dg-additional-options "-ftree-switch-conversion" } */
+
+#include <stdlib.h>
+
+int
+main (void)
+{
+  int n[1];
+
+  n[0] = 3;
+
+#pragma omp target
+  {
+    int m = n[0];
+    switch (m & 3)
+    {
+    case 0: m = 4; break;
+    case 1: m = 3; break;
+    case 2: m = 2; break;
+    default:
+      m = 1; break;
+    }
+    n[0] = m;
+  }
+
+  if (n[0] != 1)
+    abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/switch-conversion.c b/libgomp/testsuite/libgomp.c/switch-conversion.c
new file mode 100644
index 00000000000..835f54e17ac
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/switch-conversion.c
@@ -0,0 +1,36 @@
+/* PR tree-optimization/85063 */
+/* { dg-additional-options "-ftree-switch-conversion" } */
+
+#include <stdlib.h>
+
+#pragma omp declare target
+static int __attribute__((noinline)) foo (int n)
+{
+  switch (n & 3)
+    {
+    case 0: return 4;
+    case 1: return 3;
+    case 2: return 2;
+    default:
+      return 1;
+    }
+}
+#pragma omp end declare target
+
+int
+main (void)
+{
+  int n[1];
+
+  n[0] = 4;
+
+#pragma omp target
+  {
+    n[0] = foo (n[0]);
+  }
+
+  if (n[0] != 4)
+    abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/switch-conversion-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/switch-conversion-2.c
new file mode 100644
index 00000000000..8c018b80a84
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/switch-conversion-2.c
@@ -0,0 +1,31 @@
+/* PR tree-optimization/85063 */
+/* { dg-additional-options "-ftree-switch-conversion" } */
+
+#include <stdlib.h>
+
+int
+main (void)
+{
+  int n[1];
+
+  n[0] = 3;
+
+#pragma acc parallel copy(n)
+  {
+    int m = n[0];
+    switch (m & 3)
+    {
+    case 0: m = 4; break;
+    case 1: m = 3; break;
+    case 2: m = 2; break;
+    default:
+      m = 1; break;
+    }
+    n[0] = m;
+  }
+
+  if (n[0] != 1)
+    abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/switch-conversion.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/switch-conversion.c
new file mode 100644
index 00000000000..0540678d76b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/switch-conversion.c
@@ -0,0 +1,35 @@
+/* PR tree-optimization/85063 */
+/* { dg-additional-options "-ftree-switch-conversion" } */
+
+#include <stdlib.h>
+
+#pragma acc routine seq
+static int __attribute__((noinline)) foo (int n)
+{
+  switch (n & 3)
+    {
+    case 0: return 4;
+    case 1: return 3;
+    case 2: return 2;
+    default:
+      return 1;
+    }
+}
+
+int
+main (void)
+{
+  int n[1];
+
+  n[0] = 4;
+
+#pragma acc parallel copy(n)
+  {
+    n[0] = foo (n[0]);
+  }
+
+  if (n[0] != 4)
+    abort ();
+
+  return 0;
+}