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

[Tobias Burnus]

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

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