Hello Jakub, hi Honza,
On 8/31/20 5:53 PM, Jakub Jelinek wrote:
On Mon, Aug 03, 2020 at 05:37:40PM +0200, Tobias Burnus wrote:
It turned out that the omp_discover_declare_target_tgt_fn_r
discovered all nodes – but as it tagged the C++ alias nodes
and not the streamed-out nodes, no device function was created
...
if (node != NULL)
{
+ if (node->cpp_implicit_alias)
+ {
+ node = node->get_alias_target ();
+ if (!omp_declare_target_fn_p (node->decl))
+ ((vec<tree> *) data)->safe_push (node->decl);
+ DECL_ATTRIBUTES (node->decl)
+ = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (node->decl));
+ }
node->offloadable = 1;
I don't see what is special on cpp_implicit_alias here, compared to any
other aliases.
So, I wonder if the code shouldn't do:
...
Granted that cpp_implicit_alias is not special; however,
for the alias attribute, the name is different and, thus,
the decl is in a different sym node.
Updated version attached. Does it seem to make sense?
Tobias
tree decl = *tp;
symtab_node *node = symtab_node::get (decl);
if (node != NULL)
{
symtab_node *anode = node->ultimate_alias_target ();
if (anode && anode != node)
{
decl = anode->decl;
gcc_assert (TREE_CODE (decl) == FUNCTION_DECL);
if (omp_declare_target_fn_p (*tp)
|| lookup_attribute ("omp declare target host",
DECL_ATTRIBUTES (decl)))
return NULL_TREE;
node = anode;
}
}
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));
if (node != NULL)
{
node->offloadable = 1;
if (ENABLE_OFFLOADING)
g->have_offload = true;
}
Otherwise, if we have say:
void foo () { }
void bar () __attribute__((alias ("foo")));
void baz () __attribute__((alias ("bar")));
int
main ()
{
#pragma omp target
baz ();
}
we won't mark foo as being declare target.
Though, maybe that is not enough and we need to mark all the aliases from
node to the ultimate alias that way (so perhaps instead iterate one
get_alias_target (if node->alias) by one and mark all the decls the way the
code marks right now (i.e. pushes those non-DECL_EXTERNAL with
DECL_SAVED_TREE for further processing and add attributes to all of them?
Jakub
-----------------
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 | 27 +++++++++----
libgomp/testsuite/libgomp.c++/pr96390.C | 49 ++++++++++++++++++++++++
libgomp/testsuite/libgomp.c-c++-common/pr96390.c | 17 ++++++++
3 files changed, 86 insertions(+), 7 deletions(-)
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 32c2485abd4..1a3e7a384a8 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -196,21 +196,34 @@ 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)
{
+ 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;
+
node->offloadable = 1;
if (ENABLE_OFFLOADING)
g->have_offload = true;
}
+ 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..7f7fc430414
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c
@@ -0,0 +1,17 @@
+/* { dg-run } */
+/* { dg-additional-options "-O0 -fdump-tree-omplower" } */
+
+int foo () { return 42; }
+int bar () __attribute__((alias ("foo")));
+int baz () __attribute__((alias ("bar")));
+
+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" } } */
