public inbox for gcc-cvs@sourceware.org help / color / mirror / Atom feed
From: Tobias Burnus <burnus@gcc.gnu.org> To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-10] OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390) Date: Thu, 17 Sep 2020 17:03:46 +0000 (GMT) [thread overview] Message-ID: <20200917170346.B98A239960FD@sourceware.org> (raw) https://gcc.gnu.org/g:ef509d1985aa53a8c0875c25ad4050ea807be10e commit ef509d1985aa53a8c0875c25ad4050ea807be10e Author: Tobias Burnus <tobias@codesourcery.com> Date: Thu Sep 17 18:45:00 2020 +0200 OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390) Submitted to mainline at https://gcc.gnu.org/pipermail/gcc-patches/2020-September/554142.html gcc/ChangeLog: PR middle-end/96390 * omp-offload.c (omp_discover_declare_target_tgt_fn_r): Handle alias nodes. libgomp/ChangeLog: PR middle-end/96390 * testsuite/libgomp.c++/pr96390.C: New test. * testsuite/libgomp.c-c++-common/pr96390.c: New test. Diff: --- gcc/ChangeLog.omp | 6 +++ gcc/omp-offload.c | 50 ++++++++++++++++++++---- libgomp/ChangeLog.omp | 6 +++ libgomp/testsuite/libgomp.c++/pr96390.C | 49 +++++++++++++++++++++++ libgomp/testsuite/libgomp.c-c++-common/pr96390.c | 26 ++++++++++++ 5 files changed, 129 insertions(+), 8 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 63ccd8a3a79..f666abe13ac 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,9 @@ +2020-09-17 Tobias Burnus <tobias@codesourcery.com> + + PR middle-end/96390 + * omp-offload.c (omp_discover_declare_target_tgt_fn_r): Handle + alias nodes. + 2020-09-17 Tobias Burnus <tobias@codesourcery.com> Backport from mainline diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index 9c0ca643a2f..72be733a38d 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -198,21 +198,55 @@ omp_declare_target_var_p (tree decl) static tree omp_discover_declare_target_tgt_fn_r (tree *tp, int *walk_subtrees, void *data) { - if (TREE_CODE (*tp) == FUNCTION_DECL - && !omp_declare_target_fn_p (*tp) - && !lookup_attribute ("omp declare target host", DECL_ATTRIBUTES (*tp))) + if (TREE_CODE (*tp) == FUNCTION_DECL) { - tree id = get_identifier ("omp declare target"); - if (!DECL_EXTERNAL (*tp) && DECL_SAVED_TREE (*tp)) - ((vec<tree> *) data)->safe_push (*tp); - DECL_ATTRIBUTES (*tp) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (*tp)); + tree decl = *tp; symtab_node *node = symtab_node::get (*tp); if (node != NULL) { - node->offloadable = 1; + /* First, find final FUNCTION_DECL; find final alias target and there + ensure alias like cpp_implicit_alias are resolved by calling + ultimate_alias_target; the latter does not resolve alias_target as + node->analyzed = 0. */ + symtab_node *orig_node = node; + while (node->alias_target) + node = symtab_node::get (node->alias_target); + node = node->ultimate_alias_target (); + decl = node->decl; + + if (omp_declare_target_fn_p (decl) + || lookup_attribute ("omp declare target host", + DECL_ATTRIBUTES (decl))) + return NULL_TREE; + if (ENABLE_OFFLOADING) g->have_offload = true; + + /* Now mark original node and all alias targets for offloading. */ + node->offloadable = 1; + if (orig_node != node) + { + tree id = get_identifier ("omp declare target"); + while (orig_node->alias_target) + { + orig_node = orig_node->ultimate_alias_target (); + DECL_ATTRIBUTES (orig_node->decl) + = tree_cons (id, NULL_TREE, + DECL_ATTRIBUTES (orig_node->decl)); + orig_node = symtab_node::get (orig_node->alias_target); + } + } } + else if (omp_declare_target_fn_p (decl) + || lookup_attribute ("omp declare target host", + DECL_ATTRIBUTES (decl))) + return NULL_TREE; + + tree id = get_identifier ("omp declare target"); + if (!DECL_EXTERNAL (decl) && DECL_SAVED_TREE (decl)) + ((vec<tree> *) data)->safe_push (decl); + DECL_ATTRIBUTES (decl) = tree_cons (id, NULL_TREE, + DECL_ATTRIBUTES (decl)); } else if (TYPE_P (*tp)) *walk_subtrees = 0; diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 82e7cc2d941..297c6055293 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,9 @@ +2020-09-17 Tobias Burnus <tobias@codesourcery.com> + + PR middle-end/96390 + * testsuite/libgomp.c++/pr96390.C: New test. + * testsuite/libgomp.c-c++-common/pr96390.c: New test. + 2020-09-17 Tobias Burnus <tobias@codesourcery.com> Backport from mainline diff --git a/libgomp/testsuite/libgomp.c++/pr96390.C b/libgomp/testsuite/libgomp.c++/pr96390.C new file mode 100644 index 00000000000..098cb103919 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/pr96390.C @@ -0,0 +1,49 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O0 -fdump-tree-omplower" } */ + +#include <cstdlib> +#include <type_traits> + +template<int Dim> struct V { + int version_called; + + template<bool B = (Dim == 0), + typename = typename std::enable_if<B>::type> + V () + { + version_called = 1; + } + + template<typename TArg0, + typename = typename std::enable_if<(std::is_same<unsigned long, + typename std::decay<TArg0>::type>::value)>::type> + V (TArg0) + { + version_called = 2; + } +}; + +template<int Dim> struct S { + V<Dim> v; +}; + +int +main () +{ + int version_set[2] = {-1, -1}; + +#pragma omp target map(from: version_set[0:2]) + { + S<0> s; + version_set[0] = s.v.version_called; + V<1> v2((unsigned long) 1); + version_set[1] = v2.version_called; + } + + if (version_set[0] != 1 || version_set[1] != 2) + abort (); + return 0; +} + +/* "3" for S<0>::S, V<0>::V<>, and V<1>::V<long unsigned int>: */ +/* { dg-final { scan-tree-dump-times "__attribute__..omp declare target" 3 "omplower" } } */ diff --git a/libgomp/testsuite/libgomp.c-c++-common/pr96390.c b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c new file mode 100644 index 00000000000..3857bf3348d --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c @@ -0,0 +1,26 @@ +/* { dg-run } */ +/* { dg-additional-options "-O0 -fdump-tree-omplower" } */ + +#ifdef __cplusplus +extern "C" { +#endif + +int foo () { return 42; } +int bar () __attribute__((alias ("foo"))); +int baz () __attribute__((alias ("bar"))); + +#ifdef __cplusplus +} +#endif + + +int +main () +{ + int n; + #pragma omp target map(from:n) + n = baz (); + if (n != 42) + __builtin_abort (); +} +/* { dg-final { scan-tree-dump-times "__attribute__..omp declare target" 1 "omplower" } } */
next reply other threads:[~2020-09-17 17:03 UTC|newest] Thread overview: 2+ messages / expand[flat|nested] mbox.gz Atom feed top 2020-09-17 17:03 Tobias Burnus [this message] 2020-09-28 18:28 Tobias Burnus
Reply instructions: You may reply publicly to this message via plain-text email using any one of the following methods: * Save the following mbox file, import it into your mail client, and reply-to-all from there: mbox Avoid top-posting and favor interleaved quoting: https://en.wikipedia.org/wiki/Posting_style#Interleaved_style * Reply using the --to, --cc, and --in-reply-to switches of git-send-email(1): git send-email \ --in-reply-to=20200917170346.B98A239960FD@sourceware.org \ --to=burnus@gcc.gnu.org \ --cc=gcc-cvs@gcc.gnu.org \ /path/to/YOUR_REPLY https://kernel.org/pub/software/scm/git/docs/git-send-email.html * If your mail client supports setting the In-Reply-To header via mailto: links, try the mailto: linkBe sure your reply has a Subject: header at the top and a blank line before the message body.
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).