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

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

[Tobias Burnus]

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

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.)

OK?

Tobias

-----------------
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.diff --]
[-- Type: text/x-patch, Size: 2488 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
	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;
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" } } */

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

[Tobias Burnus]

On 8/3/20 5:37 PM, 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.)
>
> OK?
>
> Tobias
>
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

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

[Tobias Burnus]

And another PING.

On 8/17/20 9:17 AM, Tobias Burnus wrote:
> On 8/3/20 5:37 PM, 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.)
>>
>> OK?
>>
>> Tobias
>>
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

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

[Jakub Jelinek]

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

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

[Tobias Burnus]

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

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

[Jakub Jelinek]

On Mon, Sep 14, 2020 at 03:25:29PM +0200, Tobias Burnus wrote:
> Updated version attached. Does it seem to make sense?

I think you want Honza on this primarily, I'm always lost in the cgraph
alias code.

> --- 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 ();

I think the above is either you walk the aliases yourself, or use
ultimate_alias_target, but not both.  For the first one, e.g.
ultimate_alias_target_1 uses
      if (node->alias && node->analyzed)
        node = node->get_alias_target ();
in the loop.

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?
We don't really need to push the aliases into ((vec<tree> *) data),
that is only for function definitions (what will needs to be scanned for
further references), but I wonder if the aliases that are needed shouldn't
be marked node->offloadable and have "omp declare target" attribute added
for them too.

Perhaps use ultimate_alias_target (); first, handle the ultimate alias,
but keep the old node around, and if different from the ultimate alias,
do the node = node->get_alias_target () loop until you reach the ultimate
alias and for each add "omp declare target" and set node->offloadable
if not already there?

> +	  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;

	Jakub

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

[Tobias Burnus]

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

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

[-- Attachment #2: omp-tmpl-v3.diff --]
[-- Type: text/x-patch, Size: 5157 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                                | 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" } } */

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

[Tobias Burnus]

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

Hi Honza & Jakub,

@Honza: please look at the decl alias handling of the actual patch.

a minor update – testing only on gcn did turn out to be insufficient:
nvptx does not support alias, cf. PR 97102 + 97106; hence xfailed.

I also had a typo in one 'dg-do run' but as -O0 is set explicitly,
'dg-do run' does not make sense. (no dg-do = run once, dg-do run =
run with multiple options, dg-do compile = compile only).
Hence, I have now removed the dg-do line.

On 9/17/20 1:15 AM, Tobias Burnus wrote:
> 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

[-- Attachment #2: omp-tmpl-v6.diff --]
[-- Type: text/x-patch, Size: 5365 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                                | 51 ++++++++++++++++++++----
 libgomp/testsuite/libgomp.c++/pr96390.C          | 49 +++++++++++++++++++++++
 libgomp/testsuite/libgomp.c-c++-common/pr96390.c | 26 ++++++++++++
 3 files changed, 118 insertions(+), 8 deletions(-)

diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 32c2485abd4..222ebff1d1e 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -196,21 +196,56 @@ 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 ();
+		  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);
+		}
+	    }
 	}
+      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..8c770ecb80c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr96390.C
@@ -0,0 +1,49 @@
+/* { dg-additional-options "-O0 -fdump-tree-omplower" } */
+/* { dg-xfail-if "PR 97106/PR 97102 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */
+
+#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..692bd730069
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c
@@ -0,0 +1,26 @@
+/* { dg-additional-options "-O0 -fdump-tree-omplower" } */
+/* { dg-xfail-if "PR 97102/PR 97106 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */
+
+#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" } } */

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

[Jakub Jelinek]

On Tue, Sep 22, 2020 at 09:11:48AM +0200, Tobias Burnus wrote:
> > Okay – or do we find more issues?

I'm afraid so.

> +	  if (omp_declare_target_fn_p (decl)
> +	      || lookup_attribute ("omp declare target host",
> +				   DECL_ATTRIBUTES (decl)))
> +	    return NULL_TREE;

I'm worried that omp_declare_target_fn_p could be true and so this would
punt, but the intermediate aliases would be marked.
Or the aliases would be marked and the ultimate alias would not.
Consider:
int v;
#pragma omp declare target to (v)
void foo (void) { v++; }
void bar (void) __attribute__((alias ("foo")));
#pragma omp declare target to (bar)
void baz (void) __attribute__((alias ("foo")));
void qux (void) {
#pragma omp target
{
  bar (); // Here the ultimate alias is not marked, so the code marks it,
	  // and adds another "omp declare target" attribute to bar,
	  // which it shouldn't.
  baz (); // At this point, foo is marked, so the code wouldn't mark
	  // baz alias as "omp declare target".
}
}

So, I think it is fine to find the ultimate alias, but the loop to mark
the intermediate aliases should be invoked regardless of how decl is or is
not marked, and should test in each step whether it should or should not be
marked.

	Jakub

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

[Tobias Burnus]

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

On 9/22/20 9:36 AM, Jakub Jelinek wrote:

> On Tue, Sep 22, 2020 at 09:11:48AM +0200, Tobias Burnus wrote:
>>> Okay – or do we find more issues?
> I'm afraid so.
We will slowly converge, hopefully ;-)
> Consider:
> int v;
> #pragma omp declare target to (v)
> void foo (void) { v++; }
> void bar (void) __attribute__((alias ("foo")));
> #pragma omp declare target to (bar)
> void baz (void) __attribute__((alias ("foo")));
> void qux (void) {

etc.  – I did not convert this into a testcase. Should I?

Do you spot something more? Or is it now fine? (It passes on gcn +
xfailed on nvptx.)

Tobias
-----------------
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-v7.diff --]
[-- Type: text/x-patch, Size: 4852 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                                | 38 +++++++++++++++---
 libgomp/testsuite/libgomp.c++/pr96390.C          | 49 ++++++++++++++++++++++++
 libgomp/testsuite/libgomp.c-c++-common/pr96390.c | 26 +++++++++++++
 3 files changed, 107 insertions(+), 6 deletions(-)

diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 32c2485abd4..a1b12bbbb8f 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -196,21 +196,47 @@ 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 decl = *tp;
       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));
       symtab_node *node = symtab_node::get (*tp);
       if (node != NULL)
 	{
+	  while (node->alias_target)
+	    {
+	      node = node->ultimate_alias_target ();
+	      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));
+		}
+	      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;
+
+      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..8c770ecb80c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr96390.C
@@ -0,0 +1,49 @@
+/* { dg-additional-options "-O0 -fdump-tree-omplower" } */
+/* { dg-xfail-if "PR 97106/PR 97102 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */
+
+#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..692bd730069
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c
@@ -0,0 +1,26 @@
+/* { dg-additional-options "-O0 -fdump-tree-omplower" } */
+/* { dg-xfail-if "PR 97102/PR 97106 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */
+
+#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" } } */

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

[Jakub Jelinek]

On Tue, Sep 22, 2020 at 04:11:19PM +0200, Tobias Burnus wrote:
> +	  while (node->alias_target)
> +	    {
> +	      node = node->ultimate_alias_target ();

At least in theory, ultimate_alias_target can look through multiple aliases.
While it might not do that most of the time because this is executed quite
early, I think we have no guarantees it will never do it.
So I'd prefer what you had in the earlier patch, i.e. do the
ultimate_alias_target call + loop to find the ultimate node, and then
in another loop go from the original node (inclusive) up to the ultimate one
(exclusive) and do what you do in this loop now.
Does that make sense?

> +	      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));
> +		}
> +	      node = symtab_node::get (node->alias_target);

	Jakub

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

[Tobias Burnus]

On 9/22/20 4:24 PM, Jakub Jelinek wrote:
> On Tue, Sep 22, 2020 at 04:11:19PM +0200, Tobias Burnus wrote:
>> +      while (node->alias_target)
>> +        {
>> +          node = node->ultimate_alias_target ();
> At least in theory, ultimate_alias_target can look through multiple aliases.

Granted. But we need to handle two things:
* target alias (such as for __attribute__(alias(...))
   for this one, I can walk 'node = node->alias_target'
   here, I need to mark all nodes on the way.
* C++ aliases
   for this one, I used node = node->ultimate_alias_target ()
   here, I only need to mark the last one as only
   that function decl is streamed out

> While it might not do that most of the time because this is executed quite
> early, I think we have no guarantees it will never do it.
> So I'd prefer what you had in the earlier patch, i.e. do the
> ultimate_alias_target call + loop to find the ultimate node, and then
> in another loop go from the original node (inclusive) up to the ultimate one
> (exclusive) and do what you do in this loop now.
> Does that make sense?

I am lost. What do I gain by running the loops twice? Initially
the idea was to return NULL_TREE but as you example showed we need
to mark all unmarked ones until to final node.

Thus, we have to mark all – even if the final one is already
'omp declare target'.

But in that case, why can't we do it in a single loop?



If we assume that there is no c++ aliasing, we could do:

           while (node->alias_target)
             {
               // see assumption above: // node = node->ultimate_alias_target ();
               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));
                 }
               node = symtab_node::get (node->alias_target);
             }
           // all node->alias_target resolved
           node = node->ultimate_alias_target ();

That would avoid the in-between calling of ultimate_alias_target() but still
calls it if there is no alias_target or for the final alias target.

Is this (really) better?

BTW: When the assumption about the ordering completely changes,
the current __attribute__(alias(…)) testcase will fail; this might
not catch all issues but at least if it completely changes.

Tobias

-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

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

[Tobias Burnus]

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

And another try :-)

This time avoiding the ultimate_alias_target completely and just using:

+             if (node->cpp_implicit_alias)
+               node = node->get_alias_target ();

OK?

On 9/22/20 5:39 PM, Tobias Burnus wrote:

> On 9/22/20 4:24 PM, Jakub Jelinek wrote:
>> On Tue, Sep 22, 2020 at 04:11:19PM +0200, Tobias Burnus wrote:
>>> +      while (node->alias_target)
>>> +        {
>>> +          node = node->ultimate_alias_target ();
>> At least in theory, ultimate_alias_target can look through multiple
>> aliases.
> ...
-----------------
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-v8.diff --]
[-- Type: text/x-patch, Size: 4912 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                                | 40 ++++++++++++++++---
 libgomp/testsuite/libgomp.c++/pr96390.C          | 49 ++++++++++++++++++++++++
 libgomp/testsuite/libgomp.c-c++-common/pr96390.c | 26 +++++++++++++
 3 files changed, 109 insertions(+), 6 deletions(-)

diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 32c2485abd4..2b7325d5b7b 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -196,21 +196,49 @@ 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 decl = *tp;
       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));
       symtab_node *node = symtab_node::get (*tp);
       if (node != NULL)
 	{
+	  while (node->alias_target)
+	    {
+	      if (node->cpp_implicit_alias)
+		node = node->get_alias_target ();
+	      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));
+		}
+	      node = symtab_node::get (node->alias_target);
+	    }
+	  if (node->cpp_implicit_alias)
+	    node = node->get_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;
+
+      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..8c770ecb80c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr96390.C
@@ -0,0 +1,49 @@
+/* { dg-additional-options "-O0 -fdump-tree-omplower" } */
+/* { dg-xfail-if "PR 97106/PR 97102 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */
+
+#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..692bd730069
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c
@@ -0,0 +1,26 @@
+/* { dg-additional-options "-O0 -fdump-tree-omplower" } */
+/* { dg-xfail-if "PR 97102/PR 97106 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */
+
+#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" } } */

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

[Jakub Jelinek]

On Wed, Sep 23, 2020 at 03:52:12PM +0200, Tobias Burnus wrote:
> +  if (TREE_CODE (*tp) == FUNCTION_DECL)
>      {
> +      tree decl = *tp;
>        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));
>        symtab_node *node = symtab_node::get (*tp);
>        if (node != NULL)
>  	{
> +	  while (node->alias_target)
> +	    {
> +	      if (node->cpp_implicit_alias)
> +		node = node->get_alias_target ();
> +	      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));
> +		}
> +	      node = symtab_node::get (node->alias_target);
> +	    }
> +	  if (node->cpp_implicit_alias)
> +	    node = node->get_alias_target ();

Almost, the problem is that node with node->cpp_implicit_alias isn't marked
then.  And, the
> +           if (node->cpp_implicit_alias)
> +             node = node->get_alias_target ();
in the while loop looks problematic, not sure if it is ever possible
to have both alias_target and cpp_implicit_alias set, but if it would be,
there is no guarantee that node->get_alias_target ()->alias_target must be
non-NULL.

What I really meant was:
> +	  while (node->alias_target)
> +	    {
> +	      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));
> +		}
> +	      node = symtab_node::get (node->alias_target);
> +	    }
> +	  cgraph_node *new_node = node->get_ultimate_target ();
> +	  if (new_node != node)
> +	    {
> +	      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 ();
> +		}
> +	    }

	Jakub

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

[Tobias Burnus]

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

On 9/23/20 4:06 PM, Jakub Jelinek wrote:

> What I really meant was:
I did now something based on this.
>> +              gcc_assert (node->alias && node->analyzed);

I believe from previous testing that node->analyzed is 0
for the testcase at hand — and, hence, ultimate_alias_target()
did not walk node->target_alias for the testcase at hand.
(→ symtab_node::ultimate_alias_target_1)
Hence, I didn't include it in the assert. But I can re-check.

Tobias

-----------------
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-v9.diff --]
[-- Type: text/x-patch, Size: 5063 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                                | 44 ++++++++++++++++++---
 libgomp/testsuite/libgomp.c++/pr96390.C          | 49 ++++++++++++++++++++++++
 libgomp/testsuite/libgomp.c-c++-common/pr96390.c | 26 +++++++++++++
 3 files changed, 113 insertions(+), 6 deletions(-)

diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 32c2485abd4..dd5e9a2c9b2 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -196,21 +196,53 @@ 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 decl = *tp;
       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));
       symtab_node *node = symtab_node::get (*tp);
       if (node != NULL)
 	{
+	  while (node->alias_target)
+	    {
+	      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));
+		}
+	      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 = node->get_alias_target ();
+	    }
 	  node->offloadable = 1;
 	  if (ENABLE_OFFLOADING)
 	    g->have_offload = true;
 	}
+      if (omp_declare_target_fn_p (decl)
+	  || lookup_attribute ("omp declare target host",
+				    DECL_ATTRIBUTES (decl)))
+	return NULL_TREE;
+
+      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..8c770ecb80c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr96390.C
@@ -0,0 +1,49 @@
+/* { dg-additional-options "-O0 -fdump-tree-omplower" } */
+/* { dg-xfail-if "PR 97106/PR 97102 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */
+
+#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..692bd730069
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c
@@ -0,0 +1,26 @@
+/* { dg-additional-options "-O0 -fdump-tree-omplower" } */
+/* { dg-xfail-if "PR 97102/PR 97106 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */
+
+#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" } } */

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

[Jakub Jelinek]

On Wed, Sep 23, 2020 at 05:45:12PM +0200, Tobias Burnus wrote:
> On 9/23/20 4:06 PM, Jakub Jelinek wrote:
> 
> > What I really meant was:
> I did now something based on this.
> > > +              gcc_assert (node->alias && node->analyzed);
> 
> I believe from previous testing that node->analyzed is 0
> for the testcase at hand — and, hence, ultimate_alias_target()

That would be surprising, because if it is not node->analyzed, then
ultimate_alias_target_1 will not change node at all.

Anyway, the patch LGTM, thanks.

	Jakub