mirror/gcc
2
0
Fork 0
mirror of git://gcc.gnu.org/git/gcc.git synced 2025-01-12 01:54:54 +08:00
Code Issues Packages Projects Releases Wiki Activity

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

Browse Source 
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.
This commit is contained in:
Jakub Jelinek 2020-10-28 10:34:29 +01:00
parent 3f39b64e57
commit 2298ca2d3e
2 changed files with 63 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

23
gcc/omp-offload.c
View File

@ -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))

42
libgomp/testsuite/libgomp.c/target-42.c Normal file
View File

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