diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index 3e9c31d2cbed..4490701147ca 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 000000000000..fc0e26538ac9 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-42.c @@ -0,0 +1,42 @@ +#include + +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; +}