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.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@272322 138bc75d-0d04-0410-961f-82ee72b054a4

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 4f75fb12e2613221da993e5bd7e729a5148dafc0..febee6df6ea287bfb2df70a8e159a9e402f78df2 100644 (file)
--- 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 3b4fdc7a5099d6e33981c81a985d3370458af142..0b25e7100cde48795a7d68f52064b541c1084984 100644 (file)
--- 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 9fc92f38e6eaa85aaa712e7f99cd61e559334f8e..ea8cd7892078ee722a011e71ca0f1d74f664c30e 100644 (file)
--- 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 a68c6be9be5dcfc2ff84d611ad4ff75838250ba9..e3fbd6573b83944cb1f56adf57afdec578a1f99e 100644 (file)
--- 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 9a1fcff1fa3f63d127c3f6705fa4c5f6b810ff94..bc8647be2137cc9ee9677b95b872b51e05bde394 100644 (file)
--- 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 (file)
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 (file)
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