diff options
| -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 | ||||
| -rw-r--r-- | libgomp/ChangeLog | 6 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c/pr90779.c | 18 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.fortran/pr90779.f90 | 12 |
7 files changed, 91 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; diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 9a1fcff1f..bc8647b 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,9 @@ +2019-06-15 Jakub Jelinek <jakub@redhat.com> + + PR middle-end/90779 + * testsuite/libgomp.c/pr90779.c: New test. + * testsuite/libgomp.fortran/pr90779.f90: New test. + 2019-06-15 Tom de Vries <tdevries@suse.de> PR tree-optimization/90009 diff --git a/libgomp/testsuite/libgomp.c/pr90779.c b/libgomp/testsuite/libgomp.c/pr90779.c new file mode 100644 index 0000000..0dd1c10 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr90779.c @@ -0,0 +1,18 @@ +/* PR middle-end/90779 */ + +extern void abort (void); + +int +main () +{ + int i, j; + for (i = 0; i < 2; ++i) + #pragma omp target map(from: j) + { + static int k = 5; + j = ++k; + } + if (j != 7) + abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.fortran/pr90779.f90 b/libgomp/testsuite/libgomp.fortran/pr90779.f90 new file mode 100644 index 0000000..a6d687a --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/pr90779.f90 @@ -0,0 +1,12 @@ +! PR middle-end/90779 + +program pr90779 + implicit none + integer :: v(4), i + + !$omp target map(from:v) + v(:) = (/ (i, i=1,4) /) + !$omp end target + + if (any (v .ne. (/ (i, i=1,4) /))) stop 1 +end program |