Re: [Patch] OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390)

[Tobias Burnus <tobias@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 3017 bytes --]

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

[-- Attachment #2: omp-tmpl-v2.diff --]
[-- Type: text/x-patch, Size: 4289 bytes --]

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" } } */