Hi Honza – some input would be really helpful!
Hi Jakub – updated version below.
On 9/16/20 12:36 PM, Jakub Jelinek wrote:
I think you want Honza on this primarily, I'm always lost
in the cgraph alias code.
(Likewise as this thread shows)
+ while (node->alias_target)
+ node = symtab_node::get (node->alias_target);
+ node = node->ultimate_alias_target ();
I think the above is either you walk the aliases yourself, or use
ultimate_alias_target, but not both.
I think we need to distinguish between:
* aliases which end up with the same symbol name
and are stored in the ref_list; example: cpp_implicit_alias.
* aliases like with the alias attribute, which is handled
via alias_target and have different names.
Just experimentally:
* The 'while (node->alias_target)' properly resolves the
attribute testcase (libgomp.c-c++-common/pr96390.c).
Here, ultimate_alias_target () does not help as
node->analyzed == 0.
* The 'node->ultimate_alias_target ()' works for the
cpp_implicit_alias case (libgomp.c++/pr96390.C).
Just looking at the alias target does not help as in this
case, alias_target == NULL.
And the second thing is, I'm not sure how the aliases behave if the
ultimate alias target is properly marked as omp declare target, but some
of the aliases are not. The offloaded code will still call the alias,
so do we somehow arrange for the aliases to be also emitted into the
offloading LTO IL?
[...] I wonder if the aliases that are needed shouldn't
be marked node->offloadable and have "omp declare target" attribute added
for them too.
Done now.
Okay – or do we find more issues?
Tobias
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander
Walter
OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390)
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.
gcc/omp-offload.c | 50 ++++++++++++++++++++----
libgomp/testsuite/libgomp.c++/pr96390.C | 49 +++++++++++++++++++++++
libgomp/testsuite/libgomp.c-c++-common/pr96390.c | 26 ++++++++++++
3 files changed, 117 insertions(+), 8 deletions(-)
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 32c2485abd4..02dec44b8ce 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -196,21 +196,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/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" } } */
