From 2298ca2d3e133945f5034065e843e2ea0f36e0bb Mon Sep 17 00:00:00 2001
From: Jakub Jelinek <jakub@redhat.com>
Date: Wed, 28 Oct 2020 10:34:29 +0100
Subject: [PATCH] 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.
---
 gcc/omp-offload.c                       | 23 ++++++++++++++++--
 libgomp/testsuite/libgomp.c/target-42.c | 42 +++++++++++++++++++++++++++++++++
 2 files changed, 63 insertions(+), 2 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.c/target-42.c

diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 3e9c31d..4490701 100644
--- 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
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;
+}
-- 
2.7.4