diff mbox series

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

Message ID 20201028094312.GW7080@tucnak
State New
Headers show
Series [committed] openmp: Implicitly discover declare target for variants of declare variant calls | expand

Commit Message

Jakub Jelinek Oct. 28, 2020, 9:43 a.m. UTC 
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  <jakub@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.



	Jakub
diff mbox series

Patch

--- 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;
+}