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
> and one got link errors if offloading devices were configured.
> (Only with -O0 as otherwise inlining happened.)
>
> (Testcase is based on a sollve_vv testcase which in turn was
> based on an LLVM bugreport.)
> 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
> cpp_implicit_alias nodes.
>
> libgomp/ChangeLog:
>
> PR middle-end/96390
> * testsuite/libgomp.c++/pr96390.C: New test.
>
> gcc/omp-offload.c | 8 ++++++
> libgomp/testsuite/libgomp.c++/pr96390.C | 49
> +++++++++++++++++++++++++++++++++
> 2 files changed, 57 insertions(+)
>
> diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
> index 32c2485abd4..4aef7dbea6c 100644
> --- a/gcc/omp-offload.c
> +++ b/gcc/omp-offload.c
> @@ -207,6 +207,14 @@ omp_discover_declare_target_tgt_fn_r (tree *tp, int
> *walk_subtrees, void *data)
> symtab_node *node = symtab_node::get (*tp);
> 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;
> if (ENABLE_OFFLOADING)
> g->have_offload = true;
Sorry for the review delay.
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:
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
