diff options
| author | Jakub Jelinek <jakub@redhat.com> | 2020-10-28 10:34:29 +0100 |
|---|---|---|
| committer | Jakub Jelinek <jakub@redhat.com> | 2020-10-28 10:36:31 +0100 |
| commit | 2298ca2d3e133945f5034065e843e2ea0f36e0bb (patch) | |
| tree | df6cf359ee3a11e786641b73d6f51872410ee4d4 | |
| parent | 3f39b64e57ab8e8f69a017e4bd20aa6dd2aec492 (diff) | |
| download | gcc-2298ca2d3e133945f5034065e843e2ea0f36e0bb.zip gcc-2298ca2d3e133945f5034065e843e2ea0f36e0bb.tar.gz gcc-2298ca2d3e133945f5034065e843e2ea0f36e0bb.tar.bz2 | |
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.
| -rw-r--r-- | gcc/omp-offload.c | 23 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c/target-42.c | 42 |
2 files changed, 63 insertions, 2 deletions
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; +} |