From: Tom de Vries Date: Mon, 26 Mar 2018 09:45:49 +0000 (+0000) Subject: Fix switch conversion in offloading functions X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=46dbeb4085e4a5492579d06641d0aae25b80b465;p=gcc.git Fix switch conversion in offloading functions 2018-03-26 Tom de Vries 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 --- diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 9526547b4d8..a8d9e76b8f5 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,13 @@ +2018-03-26 Tom de Vries + + 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. + 2018-03-24 Richard Sandiford PR tree-optimization/84005 diff --git a/gcc/omp-general.c b/gcc/omp-general.c index 3ef6ce76eb8..cabbbbc6de2 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 481a8859cc5..66f0a33c2e2 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 *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 9cbc51db55e..0abf0283c9e 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 2da7068345c..b0470ef1b5e 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/ChangeLog b/libgomp/ChangeLog index 008b93535fb..a5a5e0631b2 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,11 @@ +2018-03-26 Tom de Vries + + 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 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 + +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 + +#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 + +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 + +#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; +}