From: Jakub Jelinek <jakub@redhat.com>
Date: Wed, 28 Oct 2020 09:34:29 +0000 (+0100)
Subject: openmp: Implicitly discover declare target for variants of declare variant calls
X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=2298ca2d3e133945f5034065e843e2ea0f36e0bb;p=gcc.git

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 3e9c31d2cbe..4490701147c 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 00000000000..fc0e26538ac
--- /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;
+}