public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-10] OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390)
@ 2020-09-17 17:03 Tobias Burnus
0 siblings, 0 replies; 2+ messages in thread
From: Tobias Burnus @ 2020-09-17 17:03 UTC (permalink / raw)
To: gcc-cvs
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" } } */
^ permalink raw reply [flat|nested] 2+ messages in thread
* [gcc/devel/omp/gcc-10] OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390)
@ 2020-09-28 18:28 Tobias Burnus
0 siblings, 0 replies; 2+ messages in thread
From: Tobias Burnus @ 2020-09-28 18:28 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:845a9a257334e5d8dff45784fe111b3d50eadee6
commit 845a9a257334e5d8dff45784fe111b3d50eadee6
Author: Tobias Burnus <tobias@codesourcery.com>
Date: Mon Sep 28 19:57:50 2020 +0200
OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390)
Backport mainline version; updates commit
ef509d1985aa53a8c0875c25ad4050ea807be10e for later review-comment
changes.
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.
(cherry picked from commit 2a10a2c0689db280ee3a94164504b7196b8370f4)
Diff:
---
gcc/ChangeLog.omp | 9 ++++
gcc/omp-offload.c | 59 +++++++++++-------------
libgomp/ChangeLog.omp | 9 ++++
libgomp/testsuite/libgomp.c++/pr96390.C | 1 -
libgomp/testsuite/libgomp.c-c++-common/pr96390.c | 1 -
5 files changed, 46 insertions(+), 33 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 8ea26fd0a32..595a8958314 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,12 @@
+2020-09-28 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline
+ 2020-09-28 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>
PR middle-end/96390
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index f20287653b8..edfac72bc7c 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -201,49 +201,46 @@ omp_discover_declare_target_tgt_fn_r (tree *tp, int *walk_subtrees, void *data)
if (TREE_CODE (*tp) == FUNCTION_DECL)
{
tree decl = *tp;
+ tree id = get_identifier ("omp declare target");
symtab_node *node = symtab_node::get (*tp);
if (node != NULL)
{
- /* 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)
+ if (!omp_declare_target_fn_p (node->decl)
+ && !lookup_attribute ("omp declare target host",
+ DECL_ATTRIBUTES (node->decl)))
{
- orig_node = orig_node->ultimate_alias_target ();
- orig_node->offloadable = 1;
- DECL_ATTRIBUTES (orig_node->decl)
- = tree_cons (id, NULL_TREE,
- DECL_ATTRIBUTES (orig_node->decl));
- orig_node = symtab_node::get (orig_node->alias_target);
+ node->offloadable = 1;
+ DECL_ATTRIBUTES (node->decl)
+ = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (node->decl));
}
+ node = symtab_node::get (node->alias_target);
}
+ symtab_node *new_node = node->ultimate_alias_target ();
+ decl = new_node->decl;
+ while (node != new_node)
+ {
+ if (!omp_declare_target_fn_p (node->decl)
+ && !lookup_attribute ("omp declare target host",
+ DECL_ATTRIBUTES (node->decl)))
+ {
+ node->offloadable = 1;
+ DECL_ATTRIBUTES (node->decl)
+ = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (node->decl));
+ }
+ gcc_assert (node->alias && node->analyzed);
+ node = node->get_alias_target ();
+ }
+ node->offloadable = 1;
+ if (ENABLE_OFFLOADING)
+ g->have_offload = true;
}
- else if (omp_declare_target_fn_p (decl)
- || lookup_attribute ("omp declare target host",
+ 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,
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index d570c5566dd..88428fb412c 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,12 @@
+2020-09-28 Tobias Burnus <tobias@codesourcery.com>
+
+ Backport from mainline
+ 2020-09-28 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-18 Tobias Burnus <tobias@codesourcery.com>
* testsuite/libgomp.c-c++-common/pr96390.c: XFAIL on nvptx.
diff --git a/libgomp/testsuite/libgomp.c++/pr96390.C b/libgomp/testsuite/libgomp.c++/pr96390.C
index 55b41afe7ee..8c770ecb80c 100644
--- a/libgomp/testsuite/libgomp.c++/pr96390.C
+++ b/libgomp/testsuite/libgomp.c++/pr96390.C
@@ -1,4 +1,3 @@
-/* { dg-do run } */
/* { dg-additional-options "-O0 -fdump-tree-omplower" } */
/* { dg-xfail-if "PR 97106/PR 97102 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/pr96390.c b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c
index 616fa2f83aa..692bd730069 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/pr96390.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c
@@ -1,4 +1,3 @@
-/* { dg-do run } */
/* { dg-additional-options "-O0 -fdump-tree-omplower" } */
/* { dg-xfail-if "PR 97102/PR 97106 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */
^ permalink raw reply [flat|nested] 2+ messages in thread
end of thread, other threads:[~2020-09-28 18:28 UTC | newest]
Thread overview: 2+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2020-09-17 17:03 [gcc/devel/omp/gcc-10] OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390) Tobias Burnus
2020-09-28 18:28 Tobias Burnus
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).