openmp: Implicitly discover declare target for variants of declare variant calls

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

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

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

diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 3e9c31d2cbede772b6acaa067caee139d001f3fc..4490701147caed9c93fda9df5708dc3a5dac04b8 100644 (file)
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -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 (tree *tp, int *walk_subtrees, void *data)
        }
       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))

diff --git a/libgomp/testsuite/libgomp.c/target-42.c b/libgomp/testsuite/libgomp.c/target-42.c
new file mode 100644 (file)
index 0000000..fc0e265
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-42.c
@@ -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;
+}