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

Submitted to mainline at
https://gcc.gnu.org/pipermail/gcc-patches/2020-September/554142.html

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.

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 63ccd8a3a797d90ec38c8ca04e99ff29d4513945..f666abe13ac5ae8b02b156718c13e4fd606d28b1 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,9 @@
+2020-09-17  Tobias Burnus  <tobias@codesourcery.com>
+
+       PR middle-end/96390
+       * omp-offload.c (omp_discover_declare_target_tgt_fn_r): Handle
+       alias nodes.
+
 2020-09-17  Tobias Burnus  <tobias@codesourcery.com>
 
        Backport from mainline

diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 9c0ca643a2ff25db859ab918ac18a7ac607e29dd..72be733a38d3eae19521a01315324ab1a7015e75 100644 (file)
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -198,21 +198,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/ChangeLog.omp b/libgomp/ChangeLog.omp
index 82e7cc2d941d572a116e3c92b120bf85e947d3e3..297c605529391f8b4e7d8fca439ff1e6459dcf5e 100644 (file)
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,9 @@
+2020-09-17  Tobias Burnus  <tobias@codesourcery.com>
+
+       PR middle-end/96390
+       * testsuite/libgomp.c++/pr96390.C: New test.
+       * testsuite/libgomp.c-c++-common/pr96390.c: New test.
+
 2020-09-17  Tobias Burnus  <tobias@codesourcery.com>
 
        Backport from mainline

diff --git a/libgomp/testsuite/libgomp.c++/pr96390.C b/libgomp/testsuite/libgomp.c++/pr96390.C
new file mode 100644 (file)
index 0000000..098cb10
--- /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 (file)
index 0000000..3857bf3
--- /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" } } */