diff options
| author | Jakub Jelinek <jakub@redhat.com> | 2019-06-15 09:09:04 +0200 |
|---|---|---|
| committer | Jakub Jelinek <jakub@gcc.gnu.org> | 2019-06-15 09:09:04 +0200 |
| commit | 211b7533bff68e5dd72e7d75249f470101759d6d (patch) | |
| tree | 2ee472c92423b3830a2afa44855e5bf7fdf62803 /gcc | |
| parent | 120a01d160cd47b4276507dc8f6c1ab571a9c006 (diff) | |
| download | gcc-211b7533bff68e5dd72e7d75249f470101759d6d.zip gcc-211b7533bff68e5dd72e7d75249f470101759d6d.tar.gz gcc-211b7533bff68e5dd72e7d75249f470101759d6d.tar.bz2 | |
re PR middle-end/90779 (Fortran array initialization in offload regions)
PR middle-end/90779
* gimplify.c: Include omp-offload.h and context.h.
(gimplify_bind_expr): Add "omp declare target" attributes
to static block scope variables inside of target region or target
functions.
* c-c++-common/goacc/routine-5.c (func2): Don't expect error for
static block scope variable in #pragma acc routine.
* testsuite/libgomp.c/pr90779.c: New test.
* testsuite/libgomp.fortran/pr90779.f90: New test.
From-SVN: r272322
Diffstat (limited to 'gcc')
| -rw-r--r-- | gcc/ChangeLog | 8 | ||||
| -rw-r--r-- | gcc/gimplify.c | 50 | ||||
| -rw-r--r-- | gcc/testsuite/ChangeLog | 6 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/routine-5.c | 2 |
4 files changed, 55 insertions, 11 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 4f75fb1..febee6d 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,11 @@ +2019-06-15 Jakub Jelinek <jakub@redhat.com> + + PR middle-end/90779 + * gimplify.c: Include omp-offload.h and context.h. + (gimplify_bind_expr): Add "omp declare target" attributes + to static block scope variables inside of target region or target + functions. + 2019-06-15 Tom de Vries <tdevries@suse.de> PR tree-optimization/90009 diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 3b4fdc7..0b25e71 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -65,6 +65,8 @@ along with GCC; see the file COPYING3. If not see #include "attribs.h" #include "asan.h" #include "dbgcnt.h" +#include "omp-offload.h" +#include "context.h" /* Hash set of poisoned variables in a bind expr. */ static hash_set<tree> *asan_poisoned_variables = NULL; @@ -1323,17 +1325,45 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p) struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; /* Mark variable as local. */ - if (ctx && ctx->region_type != ORT_NONE && !DECL_EXTERNAL (t) - && (! DECL_SEEN_IN_BIND_EXPR_P (t) - || splay_tree_lookup (ctx->variables, - (splay_tree_key) t) == NULL)) + if (ctx && ctx->region_type != ORT_NONE && !DECL_EXTERNAL (t)) { - if (ctx->region_type == ORT_SIMD - && TREE_ADDRESSABLE (t) - && !TREE_STATIC (t)) - omp_add_variable (ctx, t, GOVD_PRIVATE | GOVD_SEEN); - else - omp_add_variable (ctx, t, GOVD_LOCAL | GOVD_SEEN); + if (! DECL_SEEN_IN_BIND_EXPR_P (t) + || splay_tree_lookup (ctx->variables, + (splay_tree_key) t) == NULL) + { + if (ctx->region_type == ORT_SIMD + && TREE_ADDRESSABLE (t) + && !TREE_STATIC (t)) + omp_add_variable (ctx, t, GOVD_PRIVATE | GOVD_SEEN); + else + omp_add_variable (ctx, t, GOVD_LOCAL | GOVD_SEEN); + } + /* Static locals inside of target construct or offloaded + routines need to be "omp declare target". */ + if (TREE_STATIC (t)) + for (; ctx; ctx = ctx->outer_context) + if ((ctx->region_type & ORT_TARGET) != 0) + { + if (!lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (t))) + { + tree id = get_identifier ("omp declare target"); + DECL_ATTRIBUTES (t) + = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t)); + varpool_node *node = varpool_node::get (t); + if (node) + { + node->offloadable = 1; + if (ENABLE_OFFLOADING && !DECL_EXTERNAL (t)) + { + g->have_offload = true; + if (!in_lto_p) + vec_safe_push (offload_vars, t); + } + } + } + break; + } } DECL_SEEN_IN_BIND_EXPR_P (t) = 1; diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 9fc92f3..ea8cd78 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,9 @@ +2019-06-15 Jakub Jelinek <jakub@redhat.com> + + PR middle-end/90779 + * c-c++-common/goacc/routine-5.c (func2): Don't expect error for + static block scope variable in #pragma acc routine. + 2019-06-14 Steven G. Kargl <kargl@gcc.gnu.org> * gfortran.dg/integer_exponentiation_4.f90: Update test. diff --git a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c index a68c6be..e3fbd65 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-5.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c @@ -201,7 +201,7 @@ int func2 (int a) { extern int vb4; /* { dg-error "directive for use" } */ - static int vb5; /* { dg-error "directive for use" } */ + static int vb5; vb4 = a + 1; vb5 = vb4 + 1; |