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