[committed] openmp: Implicitly discover declare target for variants of declare variant calls

[Jakub Jelinek via Gcc-patches]

Hi!

This marks all variants of declare variant also declare target if the base
functions are called directly in target regions or declare target functions.

Bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk.

2020-10-28  Jakub Jelinek  <ja...@redhat.com>

        * omp-offload.c (omp_declare_target_tgt_fn_r): Handle direct calls to
        declare variant base functions.
        * testsuite/libgomp.c/target-42.c: New test.

--- gcc/omp-offload.c.jj        2020-10-22 09:26:21.772352421 +0200
+++ gcc/omp-offload.c   2020-10-22 18:54:40.920324767 +0200
@@ -196,7 +196,26 @@ 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)
+  if (TREE_CODE (*tp) == CALL_EXPR
+      && CALL_EXPR_FN (*tp)
+      && TREE_CODE (CALL_EXPR_FN (*tp)) == ADDR_EXPR
+      && TREE_CODE (TREE_OPERAND (CALL_EXPR_FN (*tp), 0)) == FUNCTION_DECL
+      && lookup_attribute ("omp declare variant base",
+                          DECL_ATTRIBUTES (TREE_OPERAND (CALL_EXPR_FN (*tp),
+                                                         0))))
+    {
+      tree fn = TREE_OPERAND (CALL_EXPR_FN (*tp), 0);
+      for (tree attr = DECL_ATTRIBUTES (fn); attr; attr = TREE_CHAIN (attr))
+       {
+         attr = lookup_attribute ("omp declare variant base", attr);
+         if (attr == NULL_TREE)
+           break;
+         tree purpose = TREE_PURPOSE (TREE_VALUE (attr));
+         if (TREE_CODE (purpose) == FUNCTION_DECL)
+           omp_discover_declare_target_tgt_fn_r (&purpose, walk_subtrees, 
data);
+       }
+    }
+  else if (TREE_CODE (*tp) == FUNCTION_DECL)
     {
       tree decl = *tp;
       tree id = get_identifier ("omp declare target");
@@ -237,7 +256,7 @@ omp_discover_declare_target_tgt_fn_r (tr
        }
       if (omp_declare_target_fn_p (decl)
          || lookup_attribute ("omp declare target host",
-                                   DECL_ATTRIBUTES (decl)))
+                              DECL_ATTRIBUTES (decl)))
        return NULL_TREE;
 
       if (!DECL_EXTERNAL (decl) && DECL_SAVED_TREE (decl))
--- libgomp/testsuite/libgomp.c/target-42.c.jj  2020-10-22 18:48:59.679244952 
+0200
+++ libgomp/testsuite/libgomp.c/target-42.c     2020-10-22 18:52:51.348903566 
+0200
@@ -0,0 +1,42 @@
+#include <stdio.h>
+
+int
+on_nvptx (void)
+{
+  return 1;
+}
+
+int
+on_gcn (void)
+{
+  return 2;
+}
+
+#pragma omp declare variant (on_nvptx) 
match(construct={target},device={arch(nvptx)})
+#pragma omp declare variant (on_gcn) 
match(construct={target},device={arch(gcn)})
+int
+on (void)
+{
+  return 0;
+}
+
+int
+main ()
+{
+  int v;
+  #pragma omp target map(from:v)
+  v = on ();
+  switch (v)
+    {
+    default:
+      printf ("Host fallback or unknown offloading\n");
+      break;
+    case 1:
+      printf ("Offloading to NVidia PTX\n");
+      break;
+    case 2:
+      printf ("Offloading to AMD GCN\n");
+      break;
+    }
+  return 0;
+}


        Jakub