diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index e00051bd3f7f..2d5cdf671ebd 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -59,6 +59,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "gimple-pretty-print.h"
 #include "stringpool.h"
 #include "attribs.h"
+#include "omp-offload.h"
 
 /* Lowering of OMP parallel and workshare constructs proceeds in two
    phases.  The first phase scans the function looking for OMP statements
@@ -1944,16 +1945,25 @@ create_omp_child_function (omp_context *ctx, bool task_copy)
 	g->have_offload = true;
     }
 
-  if (cgraph_node::get_create (decl)->offloadable
-      && !lookup_attribute ("omp declare target",
-                           DECL_ATTRIBUTES (current_function_decl)))
+  if (cgraph_node::get_create (decl)->offloadable)
     {
       const char *target_attr = (is_gimple_omp_offloaded (ctx->stmt)
 				 ? "omp target entrypoint"
 				 : "omp declare target");
-      DECL_ATTRIBUTES (decl)
-	= tree_cons (get_identifier (target_attr),
-		     NULL_TREE, DECL_ATTRIBUTES (decl));
+      if (lookup_attribute ("omp declare target",
+			    DECL_ATTRIBUTES (current_function_decl)))
+	{
+	  if (is_gimple_omp_offloaded (ctx->stmt))
+	    DECL_ATTRIBUTES (decl)
+	      = remove_attribute ("omp declare target",
+				  copy_list (DECL_ATTRIBUTES (decl)));
+	  else
+	    target_attr = NULL;
+	}
+      if (target_attr)
+	DECL_ATTRIBUTES (decl)
+	  = tree_cons (get_identifier (target_attr),
+		       NULL_TREE, DECL_ATTRIBUTES (decl));
     }
 
   t = build_decl (DECL_SOURCE_LOCATION (decl),
@@ -12960,6 +12970,23 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 				 gimple_build_assign (TREE_VEC_ELT (t, i),
 						      clobber));
 	  }
+	else if (omp_maybe_offloaded_ctx (ctx->outer))
+	  {
+	    tree id = get_identifier ("omp declare target");
+	    tree decl = TREE_VEC_ELT (t, i);
+	    DECL_ATTRIBUTES (decl)
+	      = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl));
+	    varpool_node *node = varpool_node::get (decl);
+	    if (node)
+	      {
+		node->offloadable = 1;
+		if (ENABLE_OFFLOADING)
+		  {
+		    g->have_offload = true;
+		    vec_safe_push (offload_vars, t);
+		  }
+	      }
+	  }
 
       tree clobber = build_clobber (ctx->record_type);
       gimple_seq_add_stmt (&olist, gimple_build_assign (ctx->sender_decl,
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index e9078278382d..0320ea6ab858 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -2696,8 +2696,16 @@ pass_omp_target_link::execute (function *fun)
     {
       gimple_stmt_iterator gsi;
       for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi))
-	if (walk_gimple_stmt (&gsi, NULL, find_link_var_op, NULL))
-	  gimple_regimplify_operands (gsi_stmt (gsi), &gsi);
+	{
+	  if (gimple_call_builtin_p (gsi_stmt (gsi), BUILT_IN_GOMP_TARGET))
+	    {
+	      /* Nullify the second argument of __builtin_GOMP_target_ext.  */
+	      gimple_call_set_arg (gsi_stmt (gsi), 1, null_pointer_node);
+	      update_stmt (gsi_stmt (gsi));
+	    }
+	  if (walk_gimple_stmt (&gsi, NULL, find_link_var_op, NULL))
+	    gimple_regimplify_operands (gsi_stmt (gsi), &gsi);
+	}
     }
 
   return 0;
diff --git a/libgomp/config/gcn/target.c b/libgomp/config/gcn/target.c
index 4016b7bb2404..a93ecc90d44f 100644
--- a/libgomp/config/gcn/target.c
+++ b/libgomp/config/gcn/target.c
@@ -65,3 +65,68 @@ omp_pause_resource_all (omp_pause_resource_t kind)
 
 ialias (omp_pause_resource)
 ialias (omp_pause_resource_all)
+
+void
+GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
+		 void **hostaddrs, size_t *sizes, unsigned short *kinds,
+		 unsigned int flags, void **depend, void **args)
+{
+  (void) device;
+  (void) fn;
+  (void) mapnum;
+  (void) hostaddrs;
+  (void) sizes;
+  (void) kinds;
+  (void) flags;
+  (void) depend;
+  (void) args;
+  __builtin_unreachable ();
+}
+
+void
+GOMP_target_data_ext (int device, size_t mapnum, void **hostaddrs,
+		      size_t *sizes, unsigned short *kinds)
+{
+  (void) device;
+  (void) mapnum;
+  (void) hostaddrs;
+  (void) sizes;
+  (void) kinds;
+  __builtin_unreachable ();
+}
+
+void
+GOMP_target_end_data (void)
+{
+  __builtin_unreachable ();
+}
+
+void
+GOMP_target_update_ext (int device, size_t mapnum, void **hostaddrs,
+			size_t *sizes, unsigned short *kinds,
+			unsigned int flags, void **depend)
+{
+  (void) device;
+  (void) mapnum;
+  (void) hostaddrs;
+  (void) sizes;
+  (void) kinds;
+  (void) flags;
+  (void) depend;
+  __builtin_unreachable ();
+}
+
+void
+GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
+			     size_t *sizes, unsigned short *kinds,
+			     unsigned int flags, void **depend)
+{
+  (void) device;
+  (void) mapnum;
+  (void) hostaddrs;
+  (void) sizes;
+  (void) kinds;
+  (void) flags;
+  (void) depend;
+  __builtin_unreachable ();
+}
diff --git a/libgomp/config/nvptx/target.c b/libgomp/config/nvptx/target.c
index 141057707908..e4140e482961 100644
--- a/libgomp/config/nvptx/target.c
+++ b/libgomp/config/nvptx/target.c
@@ -65,3 +65,68 @@ omp_pause_resource_all (omp_pause_resource_t kind)
 
 ialias (omp_pause_resource)
 ialias (omp_pause_resource_all)
+
+void
+GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
+		 void **hostaddrs, size_t *sizes, unsigned short *kinds,
+		 unsigned int flags, void **depend, void **args)
+{
+  (void) device;
+  (void) fn;
+  (void) mapnum;
+  (void) hostaddrs;
+  (void) sizes;
+  (void) kinds;
+  (void) flags;
+  (void) depend;
+  (void) args;
+  __builtin_unreachable ();
+}
+
+void
+GOMP_target_data_ext (int device, size_t mapnum, void **hostaddrs,
+		      size_t *sizes, unsigned short *kinds)
+{
+  (void) device;
+  (void) mapnum;
+  (void) hostaddrs;
+  (void) sizes;
+  (void) kinds;
+  __builtin_unreachable ();
+}
+
+void
+GOMP_target_end_data (void)
+{
+  __builtin_unreachable ();
+}
+
+void
+GOMP_target_update_ext (int device, size_t mapnum, void **hostaddrs,
+			size_t *sizes, unsigned short *kinds,
+			unsigned int flags, void **depend)
+{
+  (void) device;
+  (void) mapnum;
+  (void) hostaddrs;
+  (void) sizes;
+  (void) kinds;
+  (void) flags;
+  (void) depend;
+  __builtin_unreachable ();
+}
+
+void
+GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
+			     size_t *sizes, unsigned short *kinds,
+			     unsigned int flags, void **depend)
+{
+  (void) device;
+  (void) mapnum;
+  (void) hostaddrs;
+  (void) sizes;
+  (void) kinds;
+  (void) flags;
+  (void) depend;
+  __builtin_unreachable ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/for-3.c b/libgomp/testsuite/libgomp.c-c++-common/for-3.c
index 173ce8ecc136..285f8e9bd4de 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/for-3.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/for-3.c
@@ -9,6 +9,11 @@ void abort ();
 #define M(x, y, z) O(x, y, z)
 #define O(x, y, z) x ## _ ## y ## _ ## z
 
+#define DO_PRAGMA(x) _Pragma (#x)
+#define OMPTEAMS DO_PRAGMA (omp target teams)
+#define OMPFROM(v) DO_PRAGMA (omp target update from(v))
+#define OMPTO(v) DO_PRAGMA (omp target update to(v))
+
 #pragma omp declare target
 
 #define F distribute
@@ -81,33 +86,30 @@ int
 main ()
 {
   int err = 0;
-  #pragma omp target teams reduction(|:err)
-    {
-      err |= test_d_normal ();
-      err |= test_d_ds128_normal ();
-      err |= test_ds_normal ();
-      err |= test_ds_ds128_normal ();
-      err |= test_dpf_static ();
-      err |= test_dpf_static32 ();
-      err |= test_dpf_auto ();
-      err |= test_dpf_guided32 ();
-      err |= test_dpf_runtime ();
-      err |= test_dpf_ds128_static ();
-      err |= test_dpf_ds128_static32 ();
-      err |= test_dpf_ds128_auto ();
-      err |= test_dpf_ds128_guided32 ();
-      err |= test_dpf_ds128_runtime ();
-      err |= test_dpfs_static ();
-      err |= test_dpfs_static32 ();
-      err |= test_dpfs_auto ();
-      err |= test_dpfs_guided32 ();
-      err |= test_dpfs_runtime ();
-      err |= test_dpfs_ds128_static ();
-      err |= test_dpfs_ds128_static32 ();
-      err |= test_dpfs_ds128_auto ();
-      err |= test_dpfs_ds128_guided32 ();
-      err |= test_dpfs_ds128_runtime ();
-    }
+  err |= test_d_normal ();
+  err |= test_d_ds128_normal ();
+  err |= test_ds_normal ();
+  err |= test_ds_ds128_normal ();
+  err |= test_dpf_static ();
+  err |= test_dpf_static32 ();
+  err |= test_dpf_auto ();
+  err |= test_dpf_guided32 ();
+  err |= test_dpf_runtime ();
+  err |= test_dpf_ds128_static ();
+  err |= test_dpf_ds128_static32 ();
+  err |= test_dpf_ds128_auto ();
+  err |= test_dpf_ds128_guided32 ();
+  err |= test_dpf_ds128_runtime ();
+  err |= test_dpfs_static ();
+  err |= test_dpfs_static32 ();
+  err |= test_dpfs_auto ();
+  err |= test_dpfs_guided32 ();
+  err |= test_dpfs_runtime ();
+  err |= test_dpfs_ds128_static ();
+  err |= test_dpfs_ds128_static32 ();
+  err |= test_dpfs_ds128_auto ();
+  err |= test_dpfs_ds128_guided32 ();
+  err |= test_dpfs_ds128_runtime ();
   if (err)
     abort ();
   return 0;
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-41.c b/libgomp/testsuite/libgomp.c-c++-common/target-41.c
new file mode 100644
index 000000000000..3aca19a15f94
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-41.c
@@ -0,0 +1,28 @@
+/* PR libgomp/100573 */
+
+int
+foo (int a)
+{
+  if (a == 0)
+    {
+      int c;
+      a++;
+      #pragma omp target map(tofrom:a)
+      a = foo (a);
+      #pragma omp target data map(tofrom:a)
+      c = a != 2;
+      if (c)
+	return -1;
+      #pragma omp target enter data map(to:a)
+      #pragma omp target exit data map(from:a)
+    }
+  return a + 1;
+}
+
+int
+main ()
+{
+  if (foo (0) != 3)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-42.c b/libgomp/testsuite/libgomp.c-c++-common/target-42.c
new file mode 100644
index 000000000000..a334f477e512
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-42.c
@@ -0,0 +1,26 @@
+/* PR libgomp/100573 */
+
+int
+foo (int a)
+{
+  #pragma omp target firstprivate(a)
+  if (a == 0)
+    {
+      a++;
+      #pragma omp target map(tofrom:a)		/* { dg-warning "'target' construct inside of 'target' region" } */
+      a = foo (a);
+      #pragma omp target data map(tofrom:a)	/* { dg-warning "'target data' construct inside of 'target' region" } */
+      a++;
+      #pragma omp target enter data map(to:a)	/* { dg-warning "'target enter data' construct inside of 'target' region" } */
+      #pragma omp target exit data map(from:a)	/* { dg-warning "'target exit data' construct inside of 'target' region" } */
+    }
+  return a + 1;
+}
+
+int
+main ()
+{
+  if (foo (1) != 2)
+    __builtin_abort ();
+  return 0;
+}