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
Fix switch conversion in offloading functions
2018-03-25 Tom de Vries <t...@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;
+}