[PR85063] Fix switch conversion in offloading functions

Message ID 9e02d449-0c4a-88a8-0aac-1d353d321b74@mentor.com
State New
Headers show
Series
  • [PR85063] Fix switch conversion in offloading functions
Related show

Commit Message

Tom de Vries March 26, 2018, 9:05 a.m.
Hi,

this patch fixes an ICE that occurs in lto1 after switch conversion 
triggers in an offloading function.


Consider this OpenMP test-case:
...
#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;
}
...

When compiling this at -O2, we run into an ICE in lto1.

Switch conversion introduces a static variable CSWTCH.x, and lto1 ICEs 
because CSWITCH.x ends up in a different partition than the partition 
for main._omp_fn.0 (the offloading function corresponding to the omp 
target region).

Fixed by marking the CSWTCH.x variable with atrribute "omp declare 
target" in offloading functions.

Bootstrapped and reg-tested on x86_64.
Build x86_64 with nvptx accelerator and reg-tested libgomp.

OK for stage4 or stage1?

Thanks,
- Tom

Comments

Jakub Jelinek March 26, 2018, 9:13 a.m. | #1
On Mon, Mar 26, 2018 at 11:05:52AM +0200, Tom de Vries wrote:
> OK for stage4 or stage1?


Ok for stage4, thanks.

Just a small nit below.

> --- /dev/null

> +++ b/libgomp/testsuite/libgomp.c/switch-conversion-2.c

> @@ -0,0 +1,28 @@


No /* { dg-additional-options "-ftree-switch-conversion" } */
here, but:

> --- /dev/null

> +++ b/libgomp/testsuite/libgomp.c/switch-conversion.c

> @@ -0,0 +1,35 @@

> +/* { dg-additional-options "-ftree-switch-conversion" } */


only here.  Either you don't need it at all, the default
flags are -O2 I think, or use it explicitly in all the testcases, not just 2
of them.  And, please add
/* PR tree-optimization/85063 */
as the first line of all those new testcases.

	Jakub

Patch

Fix switch conversion in offloading functions

2018-03-25  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.

---
 gcc/omp-general.c                                  | 10 +++++++
 gcc/omp-general.h                                  |  1 +
 gcc/omp-offload.c                                  |  4 +--
 gcc/tree-switch-conversion.c                       |  5 ++++
 libgomp/testsuite/libgomp.c/switch-conversion-2.c  | 28 +++++++++++++++++
 libgomp/testsuite/libgomp.c/switch-conversion.c    | 35 ++++++++++++++++++++++
 .../switch-conversion-2.c                          | 28 +++++++++++++++++
 .../libgomp.oacc-c-c++-common/switch-conversion.c  | 34 +++++++++++++++++++++
 8 files changed, 142 insertions(+), 3 deletions(-)

diff --git a/gcc/omp-general.c b/gcc/omp-general.c
index 3ef6ce7..cabbbbc 100644
--- a/gcc/omp-general.c
+++ b/gcc/omp-general.c
@@ -612,6 +612,16 @@  oacc_get_fn_attrib (tree fn)
   return lookup_attribute (OACC_FN_ATTRIB, DECL_ATTRIBUTES (fn));
 }
 
+/* Return true if FN is an OpenMP or OpenACC offloading function.  */
+
+bool
+offloading_function_p (tree fn)
+{
+  tree attrs = DECL_ATTRIBUTES (fn);
+  return (lookup_attribute ("omp declare target", attrs)
+	  || lookup_attribute ("omp target entrypoint", attrs));
+}
+
 /* Extract an oacc execution dimension from FN.  FN must be an
    offloaded function or routine that has already had its execution
    dimensions lowered to the target-specific values.  */
diff --git a/gcc/omp-general.h b/gcc/omp-general.h
index 481a885..66f0a33 100644
--- a/gcc/omp-general.h
+++ b/gcc/omp-general.h
@@ -85,6 +85,7 @@  extern void oacc_replace_fn_attrib (tree fn, tree dims);
 extern void oacc_set_fn_attrib (tree fn, tree clauses, vec<tree> *args);
 extern tree oacc_build_routine_dims (tree clauses);
 extern tree oacc_get_fn_attrib (tree fn);
+extern bool offloading_function_p (tree fn);
 extern int oacc_get_fn_dim_size (tree fn, int axis);
 extern int oacc_get_ifn_dim_arg (const gimple *stmt);
 
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 9cbc51d..0abf028 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -1967,9 +1967,7 @@  public:
   virtual bool gate (function *fun)
     {
 #ifdef ACCEL_COMPILER
-      tree attrs = DECL_ATTRIBUTES (fun->decl);
-      return lookup_attribute ("omp declare target", attrs)
-	     || lookup_attribute ("omp target entrypoint", attrs);
+      return offloading_function_p (fun->decl);
 #else
       (void) fun;
       return false;
diff --git a/gcc/tree-switch-conversion.c b/gcc/tree-switch-conversion.c
index 2da7068..b0470ef 100644
--- a/gcc/tree-switch-conversion.c
+++ b/gcc/tree-switch-conversion.c
@@ -49,6 +49,7 @@  Software Foundation, 51 Franklin Street, Fifth Floor, Boston, MA
 #include "alloc-pool.h"
 #include "target.h"
 #include "tree-into-ssa.h"
+#include "omp-general.h"
 
 /* ??? For lang_hooks.types.type_for_mode, but is there a word_mode
    type in the GIMPLE type system that is language-independent?  */
@@ -1162,6 +1163,10 @@  build_one_array (gswitch *swtch, int num, tree arr_index_type,
       TREE_CONSTANT (decl) = 1;
       TREE_READONLY (decl) = 1;
       DECL_IGNORED_P (decl) = 1;
+      if (offloading_function_p (cfun->decl))
+	DECL_ATTRIBUTES (decl)
+	  = tree_cons (get_identifier ("omp declare target"), NULL_TREE,
+		       NULL_TREE);
       varpool_node::finalize_decl (decl);
 
       fetch = build4 (ARRAY_REF, value_type, decl, tidx, NULL_TREE,
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 0000000..ac4d467
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/switch-conversion-2.c
@@ -0,0 +1,28 @@ 
+#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 0000000..1ed4264
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/switch-conversion.c
@@ -0,0 +1,35 @@ 
+/* { 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 0000000..4857059
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/switch-conversion-2.c
@@ -0,0 +1,28 @@ 
+#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 0000000..d479b81
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/switch-conversion.c
@@ -0,0 +1,34 @@ 
+/* { 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;
+}