diff options
| -rw-r--r-- | gcc/c/c-decl.c | 22 | ||||
| -rw-r--r-- | gcc/cp/decl.c | 21 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c | 22 |
3 files changed, 59 insertions, 6 deletions
diff --git a/gcc/c/c-decl.c b/gcc/c/c-decl.c index b559ed5..3b2241b 100644 --- a/gcc/c/c-decl.c +++ b/gcc/c/c-decl.c @@ -58,6 +58,9 @@ along with GCC; see the file COPYING3. If not see #include "c-family/name-hint.h" #include "c-family/known-headers.h" #include "c-family/c-spellcheck.h" +#include "context.h" /* For 'g'. */ +#include "omp-general.h" +#include "omp-offload.h" /* For offload_vars. */ #include "tree-pretty-print.h" @@ -5658,9 +5661,22 @@ finish_decl (tree decl, location_t init_loc, tree init, DECL_ATTRIBUTES (decl)) && !lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (decl))) - DECL_ATTRIBUTES (decl) - = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, DECL_ATTRIBUTES (decl)); + { + DECL_ATTRIBUTES (decl) + = tree_cons (get_identifier ("omp declare target"), + NULL_TREE, DECL_ATTRIBUTES (decl)); + symtab_node *node = symtab_node::get (decl); + if (node != NULL) + { + node->offloadable = 1; + if (ENABLE_OFFLOADING) + { + g->have_offload = true; + if (is_a <varpool_node *> (node)) + vec_safe_push (offload_vars, decl); + } + } + } } /* This is the last point we can lower alignment so give the target the diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c index 9c7f6e5..1b671ce 100644 --- a/gcc/cp/decl.c +++ b/gcc/cp/decl.c @@ -53,7 +53,9 @@ along with GCC; see the file COPYING3. If not see #include "asan.h" #include "gcc-rich-location.h" #include "langhooks.h" +#include "context.h" /* For 'g'. */ #include "omp-general.h" +#include "omp-offload.h" /* For offload_vars. */ /* Possible cases of bad specifiers type used by bad_specifiers. */ enum bad_spec_place { @@ -8176,9 +8178,22 @@ cp_finish_decl (tree decl, tree init, bool init_const_expr_p, DECL_ATTRIBUTES (decl)) && !lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (decl))) - DECL_ATTRIBUTES (decl) - = tree_cons (get_identifier ("omp declare target"), - NULL_TREE, DECL_ATTRIBUTES (decl)); + { + DECL_ATTRIBUTES (decl) + = tree_cons (get_identifier ("omp declare target"), + NULL_TREE, DECL_ATTRIBUTES (decl)); + symtab_node *node = symtab_node::get (decl); + if (node != NULL) + { + node->offloadable = 1; + if (ENABLE_OFFLOADING) + { + g->have_offload = true; + if (is_a <varpool_node *> (node)) + vec_safe_push (offload_vars, decl); + } + } + } } /* This is the last point we can lower alignment so give the target the diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c b/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c new file mode 100644 index 0000000..c5670df --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/declare_target-1.c @@ -0,0 +1,22 @@ +/* PR c++/99509 */ + +#pragma omp declare target +int data[] = {5}; +#pragma omp end declare target + +static inline int +foo (int idx) +{ + return data[idx]; +} + +int +main () +{ + int i = -1; + #pragma omp target map(from:i) + i = foo(0); + if (i != 5) + __builtin_abort (); + return 0; +} |