Merge from trunk at:

	commit 56638b9b1853666f575928f8baf17f70e4ed3517
	Author: GCC Administrator <gccadmin@gcc.gnu.org>
	Date:   Wed Jun 17 00:16:36 2020 +0000

	    Daily bump.

114 files changed, 3490 insertions, 247 deletions

diff --git a/libgomp/testsuite/Makefile.am b/libgomp/testsuite/Makefile.am
index 62b1855..655a413 100644
--- a/libgomp/testsuite/Makefile.am
+++ b/libgomp/testsuite/Makefile.am
@@ -12,6 +12,8 @@ _RUNTEST = $(shell if test -f $(top_srcdir)/../dejagnu/runtest; then \
 	     echo $(top_srcdir)/../dejagnu/runtest; else echo runtest; fi)
 RUNTESTDEFAULTFLAGS = --tool $$tool --srcdir $$srcdir
 
+EXTRA_DEJAGNU_SITE_CONFIG = libgomp-site-extra.exp
+
 # Instead of directly in ../testsuite/libgomp-test-support.exp.in, the
 # following variables have to be "routed through" this Makefile, for expansion
 # of the several (Makefile) variables used therein.
diff --git a/libgomp/testsuite/Makefile.in b/libgomp/testsuite/Makefile.in
index f0da16d..52aa6c5 100644
--- a/libgomp/testsuite/Makefile.in
+++ b/libgomp/testsuite/Makefile.in
@@ -99,17 +99,20 @@ am__aclocal_m4_deps = $(top_srcdir)/../config/acx.m4 \
 	$(top_srcdir)/../config/lthostflags.m4 \
 	$(top_srcdir)/../config/multi.m4 \
 	$(top_srcdir)/../config/override.m4 \
-	$(top_srcdir)/../config/tls.m4 $(top_srcdir)/../ltoptions.m4 \
-	$(top_srcdir)/../ltsugar.m4 $(top_srcdir)/../ltversion.m4 \
-	$(top_srcdir)/../lt~obsolete.m4 $(top_srcdir)/acinclude.m4 \
-	$(top_srcdir)/../libtool.m4 $(top_srcdir)/../config/cet.m4 \
+	$(top_srcdir)/../config/tls.m4 \
+	$(top_srcdir)/../config/toolexeclibdir.m4 \
+	$(top_srcdir)/../ltoptions.m4 $(top_srcdir)/../ltsugar.m4 \
+	$(top_srcdir)/../ltversion.m4 $(top_srcdir)/../lt~obsolete.m4 \
+	$(top_srcdir)/acinclude.m4 $(top_srcdir)/../libtool.m4 \
+	$(top_srcdir)/../config/cet.m4 \
 	$(top_srcdir)/plugin/configfrag.ac $(top_srcdir)/configure.ac
 am__configure_deps = $(am__aclocal_m4_deps) $(CONFIGURE_DEPENDENCIES) \
 	$(ACLOCAL_M4)
 DIST_COMMON = $(srcdir)/Makefile.am
 mkinstalldirs = $(SHELL) $(top_srcdir)/../mkinstalldirs
 CONFIG_HEADER = $(top_builddir)/config.h
-CONFIG_CLEAN_FILES = libgomp-test-support.pt.exp
+CONFIG_CLEAN_FILES = libgomp-test-support.pt.exp \
+	libgomp-site-extra.exp
 CONFIG_CLEAN_VPATH_FILES =
 AM_V_P = $(am__v_P_@AM_V@)
 am__v_P_ = $(am__v_P_@AM_DEFAULT_V@)
@@ -308,6 +311,7 @@ _RUNTEST = $(shell if test -f $(top_srcdir)/../dejagnu/runtest; then \
 	     echo $(top_srcdir)/../dejagnu/runtest; else echo runtest; fi)
 
 RUNTESTDEFAULTFLAGS = --tool $$tool --srcdir $$srcdir
+EXTRA_DEJAGNU_SITE_CONFIG = libgomp-site-extra.exp
 all: all-am
 
 .SUFFIXES:
@@ -342,6 +346,8 @@ $(ACLOCAL_M4): @MAINTAINER_MODE_TRUE@ $(am__aclocal_m4_deps)
 $(am__aclocal_m4_deps):
 libgomp-test-support.pt.exp: $(top_builddir)/config.status $(srcdir)/libgomp-test-support.exp.in
 	cd $(top_builddir) && $(SHELL) ./config.status $(subdir)/$@
+libgomp-site-extra.exp: $(top_builddir)/config.status $(srcdir)/libgomp-site-extra.exp.in
+	cd $(top_builddir) && $(SHELL) ./config.status $(subdir)/$@
 
 mostlyclean-libtool:
 	-rm -f *.lo
diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index f52ed71..ee5f0e5 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -319,7 +319,7 @@ proc libgomp_option_proc { option } {
 proc offload_target_to_openacc_device_type { offload_target } {
     switch -glob $offload_target {
 	amdgcn* {
-	    return "gcn"
+	    return "radeon"
 	}
 	disable {
 	    return "host"
@@ -346,11 +346,11 @@ proc check_effective_target_offload_target_nvptx { } {
     # files; in particular, '-foffload', 'libgomp.oacc-*/*.exp'), which don't
     # get passed on to 'check_effective_target_*' functions.  (Not caching the
     # result due to that.)
-    set options [current_compiler_flags]
+    set options [list "additional_flags=[concat "-v" [current_compiler_flags]]"]
     # Instead of inspecting command-line options, look what the compiler driver
     # decides.  This is somewhat modelled after
     # 'gcc/testsuite/lib/target-supports.exp:check_configured_with'.
-    set gcc_output [libgomp_target_compile "-v $options" "" "none" ""]
+    set gcc_output [libgomp_target_compile "" "" "none" $options]
     if [regexp "(?n)^OFFLOAD_TARGET_NAMES=(.*)" $gcc_output dummy offload_targets] {
 	verbose "compiling for offload targets: $offload_targets"
 	return [string match "*:nvptx*:*" ":$offload_targets:"]
@@ -483,22 +483,22 @@ proc check_effective_target_hsa_offloading_selected {} {
     }]
 }
 
-# Return 1 if at least one AMD GCN board is present.
+# Return 1 if at least one AMD GPU is accessible.
 
-proc check_effective_target_openacc_amdgcn_accel_present { } {
-    return [check_runtime openacc_amdgcn_accel_present {
+proc check_effective_target_openacc_radeon_accel_present { } {
+    return [check_runtime openacc_radeon_accel_present {
 	#include <openacc.h>
 	int main () {
-	return !(acc_get_num_devices (acc_device_gcn) > 0);
+	return !(acc_get_num_devices (acc_device_radeon) > 0);
 	}
     } "" ]
 }
 
-# Return 1 if at least one AMD GCN board is present, and the AMD GCN device
-# type is selected by default.
+# Return 1 if at least one AMD GPU is accessible, and the OpenACC 'radeon'
+# device type is selected.
 
-proc check_effective_target_openacc_amdgcn_accel_selected { } {
-    if { ![check_effective_target_openacc_amdgcn_accel_present] } {
+proc check_effective_target_openacc_radeon_accel_selected { } {
+    if { ![check_effective_target_openacc_radeon_accel_present] } {
 	return 0;
     }
     global offload_target
diff --git a/libgomp/testsuite/libgomp-site-extra.exp.in b/libgomp/testsuite/libgomp-site-extra.exp.in
new file mode 100644
index 0000000..c0d2666
--- /dev/null
+++ b/libgomp/testsuite/libgomp-site-extra.exp.in
@@ -0,0 +1 @@
+set GCC_UNDER_TEST {@CC@}
diff --git a/libgomp/testsuite/libgomp-test-support.exp.in b/libgomp/testsuite/libgomp-test-support.exp.in
index 6ec10c7..98fb442 100644
--- a/libgomp/testsuite/libgomp-test-support.exp.in
+++ b/libgomp/testsuite/libgomp-test-support.exp.in
@@ -1,5 +1,3 @@
-set GCC_UNDER_TEST {@CC@}
-
 set cuda_driver_include "@CUDA_DRIVER_INCLUDE@"
 set cuda_driver_lib "@CUDA_DRIVER_LIB@"
 set hsa_runtime_lib "@HSA_RUNTIME_LIB@"
diff --git a/libgomp/testsuite/libgomp.c++/pr93931.C b/libgomp/testsuite/libgomp.c++/pr93931.C
new file mode 100644
index 0000000..4d4232e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/pr93931.C
@@ -0,0 +1,120 @@
+// PR c++/93931
+// { dg-do run }
+// { dg-options "-O2 -std=c++14" }
+
+extern "C" void abort ();
+
+void
+sink (int &x)
+{
+  int *volatile p;
+  p = &x;
+  (*p)++;
+}
+
+int
+foo ()
+{
+  int r = 0;
+  [&r] () {
+#pragma omp parallel for reduction(+ : r)
+    for (int i = 0; i < 1024; ++i)
+      r += i;
+  } ();
+  return r;
+}
+
+int
+bar ()
+{
+  int l = 0;
+  [&l] () {
+#pragma omp parallel for lastprivate (l)
+    for (int i = 0; i < 1024; ++i)
+      l = i;
+  } ();
+  return l;
+}
+
+void
+baz ()
+{
+  int f = 18;
+  [&f] () {
+#pragma omp parallel for firstprivate (f)
+    for (int i = 0; i < 1024; ++i)
+      {
+	sink (f);
+	f += 3;
+	sink (f);
+	if (f != 23)
+	  abort ();
+	sink (f);
+	f -= 7;
+	sink (f);
+      }
+  } ();
+  if (f != 18)
+    abort ();
+}
+
+int
+qux ()
+{
+  int r = 0;
+  [&] () {
+#pragma omp parallel for reduction(+ : r)
+    for (int i = 0; i < 1024; ++i)
+      r += i;
+  } ();
+  return r;
+}
+
+int
+corge ()
+{
+  int l = 0;
+  [&] () {
+#pragma omp parallel for lastprivate (l)
+    for (int i = 0; i < 1024; ++i)
+      l = i;
+  } ();
+  return l;
+}
+
+void
+garply ()
+{
+  int f = 18;
+  [&] () {
+#pragma omp parallel for firstprivate (f)
+    for (int i = 0; i < 1024; ++i)
+      {
+	sink (f);
+	f += 3;
+	sink (f);
+	if (f != 23)
+	  abort ();
+	sink (f);
+	f -= 7;
+	sink (f);
+      }
+  } ();
+  if (f != 18)
+    abort ();
+}
+
+int
+main ()
+{
+  if (foo () != 1024 * 1023 / 2)
+    abort ();
+  if (bar () != 1023)
+    abort ();
+  baz ();
+  if (qux () != 1024 * 1023 / 2)
+    abort ();
+  if (corge () != 1023)
+    abort ();
+  garply ();
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/alloc-1.c b/libgomp/testsuite/libgomp.c-c++-common/alloc-1.c
new file mode 100644
index 0000000..9259a9c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/alloc-1.c
@@ -0,0 +1,157 @@
+#include <omp.h>
+#include <stdint.h>
+#include <stdlib.h>
+
+const omp_alloctrait_t traits2[]
+= { { omp_atk_alignment, 16 },
+    { omp_atk_sync_hint, omp_atv_default },
+    { omp_atk_access, omp_atv_default },
+    { omp_atk_pool_size, 1024 },
+    { omp_atk_fallback, omp_atv_default_mem_fb },
+    { omp_atk_partition, omp_atv_environment } };
+omp_alloctrait_t traits3[]
+= { { omp_atk_sync_hint, omp_atv_uncontended },
+    { omp_atk_alignment, 32 },
+    { omp_atk_access, omp_atv_all },
+    { omp_atk_pool_size, 512 },
+    { omp_atk_fallback, omp_atv_allocator_fb },
+    { omp_atk_fb_data, 0 },
+    { omp_atk_partition, omp_atv_default } };
+const omp_alloctrait_t traits4[]
+= { { omp_atk_alignment, 128 },
+    { omp_atk_pool_size, 1024 },
+    { omp_atk_fallback, omp_atv_null_fb } };
+
+int
+main ()
+{
+  int *volatile p = (int *) omp_alloc (3 * sizeof (int), omp_default_mem_alloc);
+  int *volatile q;
+  int *volatile r;
+  omp_alloctrait_t traits[3]
+    = { { omp_atk_alignment, 64 },
+	{ omp_atk_fallback, omp_atv_null_fb },
+	{ omp_atk_pool_size, 4096 } };
+  omp_allocator_handle_t a, a2;
+
+  if ((((uintptr_t) p) % __alignof (int)) != 0)
+    abort ();
+  p[0] = 1;
+  p[1] = 2;
+  p[2] = 3;
+  omp_free (p, omp_default_mem_alloc);
+  p = (int *) omp_alloc (2 * sizeof (int), omp_default_mem_alloc);
+  if ((((uintptr_t) p) % __alignof (int)) != 0)
+    abort ();
+  p[0] = 1;
+  p[1] = 2;
+  omp_free (p, omp_null_allocator);
+  omp_set_default_allocator (omp_default_mem_alloc);
+  p = (int *) omp_alloc (sizeof (int), omp_null_allocator);
+  if ((((uintptr_t) p) % __alignof (int)) != 0)
+    abort ();
+  p[0] = 3;
+  omp_free (p, omp_get_default_allocator ());
+
+  a = omp_init_allocator (omp_default_mem_space, 3, traits);
+  if (a == omp_null_allocator)
+    abort ();
+  p = (int *) omp_alloc (3072, a);
+  if ((((uintptr_t) p) % 64) != 0)
+    abort ();
+  p[0] = 1;
+  p[3071 / sizeof (int)] = 2;
+  if (omp_alloc (3072, a) != NULL)
+    abort ();
+  omp_free (p, a);
+  p = (int *) omp_alloc (3072, a);
+  p[0] = 3;
+  p[3071 / sizeof (int)] = 4;
+  omp_free (p, omp_null_allocator);
+  omp_set_default_allocator (a);
+  if (omp_get_default_allocator () != a)
+    abort ();
+  p = (int *) omp_alloc (3072, omp_null_allocator);
+  if (omp_alloc (3072, omp_null_allocator) != NULL)
+    abort ();
+  omp_free (p, a);
+  omp_destroy_allocator (a);
+
+  a = omp_init_allocator (omp_default_mem_space,
+			  sizeof (traits2) / sizeof (traits2[0]),
+			  traits2);
+  if (a == omp_null_allocator)
+    abort ();
+  if (traits3[5].key != omp_atk_fb_data)
+    abort ();
+  traits3[5].value = (uintptr_t) a;
+  a2 = omp_init_allocator (omp_default_mem_space,
+			   sizeof (traits3) / sizeof (traits3[0]),
+			   traits3);
+  if (a2 == omp_null_allocator)
+    abort ();
+  p = (int *) omp_alloc (420, a2);
+  if ((((uintptr_t) p) % 32) != 0)
+    abort ();
+  p[0] = 5;
+  p[419 / sizeof (int)] = 6;
+  q = (int *) omp_alloc (768, a2);
+  if ((((uintptr_t) q) % 16) != 0)
+    abort ();
+  q[0] = 7;
+  q[767 / sizeof (int)] = 8;
+  r = (int *) omp_alloc (512, a2);
+  if ((((uintptr_t) r) % __alignof (int)) != 0)
+    abort ();
+  r[0] = 9;
+  r[511 / sizeof (int)] = 10;
+  omp_free (p, omp_null_allocator);
+  omp_free (q, a2);
+  omp_free (r, omp_null_allocator);
+  omp_destroy_allocator (a2);
+  omp_destroy_allocator (a);
+
+  a = omp_init_allocator (omp_default_mem_space,
+			  sizeof (traits4) / sizeof (traits4[0]),
+			  traits4);
+  if (a == omp_null_allocator)
+    abort ();
+  if (traits3[5].key != omp_atk_fb_data)
+    abort ();
+  traits3[5].value = (uintptr_t) a;
+  a2 = omp_init_allocator (omp_default_mem_space,
+			   sizeof (traits3) / sizeof (traits3[0]),
+			   traits3);
+  if (a2 == omp_null_allocator)
+    abort ();
+  omp_set_default_allocator (a2);
+#ifdef __cplusplus
+  p = static_cast <int *> (omp_alloc (420));
+#else
+  p = (int *) omp_alloc (420, omp_null_allocator);
+#endif
+  if ((((uintptr_t) p) % 32) != 0)
+    abort ();
+  p[0] = 5;
+  p[419 / sizeof (int)] = 6;
+  q = (int *) omp_alloc (768, omp_null_allocator);
+  if ((((uintptr_t) q) % 128) != 0)
+    abort ();
+  q[0] = 7;
+  q[767 / sizeof (int)] = 8;
+  if (omp_alloc (768, omp_null_allocator) != NULL)
+    abort ();
+#ifdef __cplusplus
+  omp_free (p);
+  omp_free (q);
+  omp_free (NULL);
+#else
+  omp_free (p, omp_null_allocator);
+  omp_free (q, omp_null_allocator);
+  omp_free (NULL, omp_null_allocator);
+#endif
+  omp_free (NULL, omp_null_allocator);
+  omp_destroy_allocator (a2);
+  omp_destroy_allocator (a);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/alloc-2.c b/libgomp/testsuite/libgomp.c-c++-common/alloc-2.c
new file mode 100644
index 0000000..ee53958
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/alloc-2.c
@@ -0,0 +1,46 @@
+#include <omp.h>
+#include <stdint.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  omp_alloctrait_t traits[3]
+    = { { omp_atk_alignment, 64 },
+	{ omp_atk_fallback, omp_atv_null_fb },
+	{ omp_atk_pool_size, 4096 } };
+  omp_allocator_handle_t a
+    = omp_init_allocator (omp_default_mem_space, 3, traits);
+  if (a == omp_null_allocator)
+    abort ();
+
+  #pragma omp parallel num_threads(4)
+  {
+    int n = omp_get_thread_num ();
+    double *volatile p, *volatile q;
+    omp_set_default_allocator ((n & 1) ? a : omp_default_mem_alloc);
+    p = (double *) omp_alloc (1696, omp_null_allocator);
+    if (p == NULL)
+      abort ();
+    p[0] = 1.0;
+    p[1695 / sizeof (double *)] = 2.0;
+    #pragma omp barrier
+    omp_set_default_allocator ((n & 1) ? omp_default_mem_alloc : a);
+    q = (double *) omp_alloc (1696, omp_null_allocator);
+    if (n & 1)
+      {
+	if (q == NULL)
+	  abort ();
+	q[0] = 3.0;
+	q[1695 / sizeof (double *)] = 4.0;
+      }
+    else if (q != NULL)
+      abort ();
+    #pragma omp barrier
+    omp_free (p, omp_null_allocator);
+    omp_free (q, omp_null_allocator);
+    omp_set_default_allocator (omp_default_mem_alloc);
+  }
+  omp_destroy_allocator (a);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/alloc-3.c b/libgomp/testsuite/libgomp.c-c++-common/alloc-3.c
new file mode 100644
index 0000000..a30cdc0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/alloc-3.c
@@ -0,0 +1,28 @@
+/* { dg-set-target-env-var OMP_ALLOCATOR "omp_cgroup_mem_alloc" } */
+/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */
+
+#include <string.h>
+#include <stdlib.h>
+#include <omp.h>
+
+int
+main ()
+{
+  const char *p = getenv ("OMP_ALLOCATOR");
+  if (p && strcmp (p, "omp_cgroup_mem_alloc") == 0)
+    {
+      if (omp_get_default_allocator () != omp_cgroup_mem_alloc)
+	abort ();
+      #pragma omp parallel num_threads (2)
+      {
+	if (omp_get_default_allocator () != omp_cgroup_mem_alloc)
+	  abort ();
+	#pragma omp parallel num_threads (2)
+	{
+	  if (omp_get_default_allocator () != omp_cgroup_mem_alloc)
+	    abort ();
+	}
+      }
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/alloc-4.c b/libgomp/testsuite/libgomp.c-c++-common/alloc-4.c
new file mode 100644
index 0000000..841e1bc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/alloc-4.c
@@ -0,0 +1,25 @@
+#include <omp.h>
+#include <stdlib.h>
+
+const omp_alloctrait_t traits[]
+= { { omp_atk_pool_size, 1 },
+    { omp_atk_fallback, omp_atv_abort_fb } };
+
+int
+main ()
+{
+  omp_allocator_handle_t a;
+
+  if (omp_alloc (0, omp_null_allocator) != NULL)
+    abort ();
+  a = omp_init_allocator (omp_default_mem_space, 2, traits);
+  if (a != omp_null_allocator)
+    {
+      if (omp_alloc (0, a) != NULL
+	  || omp_alloc (0, a) != NULL
+	  || omp_alloc (0, a) != NULL)
+	abort ();
+      omp_destroy_allocator (a);
+    }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/pr93515.c b/libgomp/testsuite/libgomp.c-c++-common/pr93515.c
new file mode 100644
index 0000000..8a69088
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/pr93515.c
@@ -0,0 +1,36 @@
+/* PR libgomp/93515 */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  int i;
+  int a = 42;
+#pragma omp target teams distribute parallel for defaultmap(tofrom: scalar)
+  for (i = 0; i < 64; ++i)
+    if (omp_get_team_num () == 0)
+      if (omp_get_thread_num () == 0)
+	a = 142;
+  if (a != 142)
+    __builtin_abort ();
+  a = 42;
+#pragma omp target parallel for defaultmap(tofrom: scalar)
+  for (i = 0; i < 64; ++i)
+    if (omp_get_thread_num () == 0)
+      a = 143;
+  if (a != 143)
+    __builtin_abort ();
+  a = 42;
+#pragma omp target firstprivate(a)
+  {
+    #pragma omp parallel for
+    for (i = 0; i < 64; ++i)
+      if (omp_get_thread_num () == 0)
+	a = 144;
+    if (a != 144)
+      abort ();
+  }
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-40.c b/libgomp/testsuite/libgomp.c-c++-common/target-40.c
new file mode 100644
index 0000000..22bbdd9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-40.c
@@ -0,0 +1,51 @@
+/* { dg-do run } */
+/* { dg-options "-O0" } */
+
+extern
+#ifdef __cplusplus
+"C"
+#endif
+void abort (void);
+volatile int v;
+#pragma omp declare target to (v)
+typedef void (*fnp1) (void);
+typedef fnp1 (*fnp2) (void);
+void f1 (void) { v++; }
+void f2 (void) { v += 4; }
+void f3 (void) { v += 16; f1 (); }
+fnp1 f4 (void) { v += 64; return f2; }
+int a = 1;
+int *b = &a;
+int **c = &b;
+fnp2 f5 (void) { f3 (); return f4; }
+#pragma omp declare target to (c)
+
+int
+main ()
+{
+  int err = 0;
+  #pragma omp target map(from:err)
+  {
+    volatile int xa;
+    int *volatile xb;
+    int **volatile xc;
+    fnp2 xd;
+    fnp1 xe;
+    err = 0;
+    xa = a;
+    err |= xa != 1;
+    xb = b;
+    err |= xb != &a;
+    xc = c;
+    err |= xc != &b;
+    xd = f5 ();
+    err |= v != 17;
+    xe = xd ();
+    err |= v != 81;
+    xe ();
+    err |= v != 85;
+  }
+  if (err)
+    abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/pr93566.c b/libgomp/testsuite/libgomp.c/pr93566.c
new file mode 100644
index 0000000..3334bd57
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/pr93566.c
@@ -0,0 +1,113 @@
+/* PR middle-end/93566 */
+/* { dg-additional-options "-std=c99" } */
+
+extern void abort (void);
+
+void
+foo (int *x)
+{
+  void nest (void) {
+    #pragma omp parallel for reduction(+:x[:10])
+    for (int i = 0; i < 1024; i++)
+      for (int j = 0; j < 10; j++)
+	x[j] += j * i;
+  }
+  nest ();
+  for (int i = 0; i < 10; i++)
+    if (x[i] != 1023 * 1024 / 2 * i)
+      abort ();
+}
+
+void
+bar (void)
+{
+  int x[10] = {};
+  void nest (void) {
+    #pragma omp parallel for reduction(+:x[:10])
+    for (int i = 0; i < 1024; i++)
+      for (int j = 0; j < 10; j++)
+	x[j] += j * i;
+  }
+  nest ();
+  for (int i = 0; i < 10; i++)
+    if (x[i] != 1023 * 1024 / 2 * i)
+      abort ();
+}
+
+void
+baz (void)
+{
+  int x[10] = {};
+  void nest (void) {
+    #pragma omp parallel for reduction(+:x[2:5])
+    for (int i = 0; i < 1024; i++)
+      for (int j = 2; j < 7; j++)
+	x[j] += j * i;
+  }
+  nest ();
+  for (int i = 2; i < 7; i++)
+    if (x[i] != 1023 * 1024 / 2 * i)
+      abort ();
+}
+
+void
+qux (int *x)
+{
+  void nest (void) { x++; }
+  nest ();
+  #pragma omp parallel for reduction(+:x[:9])
+  for (int i = 0; i < 1024; i++)
+    for (int j = 0; j < 9; j++)
+      x[j] += j * i;
+  nest ();
+  for (int i = 0; i < 9; i++)
+    if (x[i - 1] != 1023 * 1024 / 2 * i)
+      abort ();
+}
+
+void
+quux (void)
+{
+  int x[10];
+  void nest (void) { for (int i = 0; i < 10; i++) x[i] = 0; }
+  int nest2 (int i) { return x[i]; }
+  nest ();
+  #pragma omp parallel for reduction(+:x[:7])
+  for (int i = 0; i < 1024; i++)
+    for (int j = 0; j < 7; j++)
+      x[j] += j * i;
+  for (int i = 0; i < 7; i++)
+    if (nest2 (i) != 1023 * 1024 / 2 * i)
+      abort ();
+}
+
+void
+corge (void)
+{
+  int x[10];
+  void nest (void) { for (int i = 0; i < 10; i++) x[i] = 0; }
+  int nest2 (int i) { return x[i]; }
+  nest ();
+  #pragma omp parallel for reduction(+:x[2:4])
+  for (int i = 0; i < 1024; i++)
+    for (int j = 2; j < 6; j++)
+      x[j] += j * i;
+  for (int i = 2; i < 6; i++)
+    if (nest2 (i) != 1023 * 1024 / 2 * i)
+      abort ();
+}
+
+int
+main ()
+{
+  int a[10] = {};
+  foo (a);
+  bar ();
+  baz ();
+  for (int i = 0; i < 10; i++)
+    a[i] = 0;
+  qux (a);
+  quux ();
+  corge ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-38.c b/libgomp/testsuite/libgomp.c/target-38.c
new file mode 100644
index 0000000..8169972
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-38.c
@@ -0,0 +1,28 @@
+#define A(n) n##0, n##1, n##2, n##3, n##4, n##5, n##6, n##7, n##8, n##9
+#define B(n) A(n##0), A(n##1), A(n##2), A(n##3), A(n##4), A(n##5), A(n##6), A(n##7), A(n##8), A(n##9)
+
+int
+foo (int x)
+{
+  int b[] = { B(4), B(5), B(6) };
+  return b[x];
+}
+
+int v[] = { 1, 2, 3, 4, 5, 6 };
+#pragma omp declare target to (foo, v)
+
+int
+main ()
+{
+  int i = 5;
+  asm ("" : "+g" (i));
+  #pragma omp target map(tofrom:i)
+  {
+    int a[] = { B(1), B(2), B(3) };
+    asm ("" : : "m" (a) : "memory");
+    i = a[i] + foo (i) + v[i & 63];
+  }
+  if (i != 105 + 405 + 6)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-39.c b/libgomp/testsuite/libgomp.c/target-39.c
new file mode 100644
index 0000000..4442f43
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-39.c
@@ -0,0 +1,47 @@
+/* { dg-do run } */
+/* { dg-options "-O0" } */
+
+extern void abort (void);
+volatile int v;
+#pragma omp declare target to (v)
+typedef void (*fnp1) (void);
+typedef fnp1 (*fnp2) (void);
+void f1 (void) { v++; }
+void f2 (void) { v += 4; }
+void f3 (void) { v += 16; f1 (); }
+fnp1 f4 (void) { v += 64; return f2; }
+int a = 1;
+int *b = &a;
+int **c = &b;
+fnp2 f5 (void) { f3 (); return f4; }
+#pragma omp declare target to (c, f5)
+
+int
+main ()
+{
+  int err = 0;
+  #pragma omp target map(from:err)
+  {
+    volatile int xa;
+    int *volatile xb;
+    int **volatile xc;
+    fnp2 xd;
+    fnp1 xe;
+    err = 0;
+    xa = a;
+    err |= xa != 1;
+    xb = b;
+    err |= xb != &a;
+    xc = c;
+    err |= xc != &b;
+    xd = f5 ();
+    err |= v != 17;
+    xe = xd ();
+    err |= v != 81;
+    xe ();
+    err |= v != 85;
+  }
+  if (err)
+    abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/async_io_9.f90 b/libgomp/testsuite/libgomp.fortran/async_io_9.f90
new file mode 100644
index 0000000..2dc111c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/async_io_9.f90
@@ -0,0 +1,20 @@
+! { dg-do run }
+! PR 95191 - this used to hang.
+! Original test case by Bill Long.
+program test
+  real a(10000)
+  integer my_id
+  integer bad_id
+  integer :: iostat
+  character (len=100) :: iomsg
+  data my_id /1/
+  data bad_id /2/
+  a = 1.
+  open (unit=10, file='test.dat', form='unformatted', &
+       &                asynchronous='yes')
+  write (unit=10, asynchronous='yes', id=my_id) a
+  iomsg = ""
+  wait (unit=10, id=bad_id, iostat=iostat, iomsg=iomsg)
+  if (iostat == 0 .or. iomsg /= "Bad ID in WAIT statement") stop 1
+  close (unit=10, status='delete')
+end program test
diff --git a/libgomp/testsuite/libgomp.fortran/close_errors_1.f90 b/libgomp/testsuite/libgomp.fortran/close_errors_1.f90
new file mode 100644
index 0000000..6edb7da
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/close_errors_1.f90
@@ -0,0 +1,19 @@
+! { dg-do run }
+! PR 95115 - this used to hang with -pthread.  Original test case by
+! Bill Long.
+
+program test
+  character(len=16) my_status
+  character(len=1000) :: iomsg
+  open (unit=10, file='test.dat')
+  print *,42
+  write (10, *) 'weird'
+  rewind (10)
+  read (10, *) my_status
+  close (10)
+  open (unit=10, file='test.dat')
+  close (unit=10, status=my_status, iostat=ios, iomsg=iomsg)
+  if (ios == 0) stop 1
+  if (iomsg /= "Bad STATUS parameter in CLOSE statement") stop 2
+  close (10, status='delete')
+end program test
diff --git a/libgomp/testsuite/libgomp.fortran/pr66199-3.f90 b/libgomp/testsuite/libgomp.fortran/pr66199-3.f90
new file mode 100644
index 0000000..7c596dc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr66199-3.f90
@@ -0,0 +1,53 @@
+! { dg-do run }
+!
+! PR fortran/94690
+! PR middle-end/66199 
+
+module m
+  integer u(0:1024-1), v(0:1024-1), w(0:1024-1)
+contains
+
+integer(8) function f1 (a, b)
+  implicit none
+  integer, value :: a, b
+  integer(8) :: d
+  !$omp parallel do lastprivate (d) default(none) firstprivate (a, b) shared(u, v, w)
+  do d = a, b-1
+    u(d) = v(d) + w(d)
+  end do
+  f1 = d
+end
+
+integer(8) function f2 (a, b, c)
+  implicit none
+  integer, value :: a, b, c
+  integer(8) :: d, e
+  !$omp parallel do lastprivate (d) default(none) firstprivate (a, b) shared(u, v, w) linear(c:5) lastprivate(e)
+  do d = a, b-1
+      u(d) = v(d) + w(d)
+      c = c + 5
+      e = c
+  end do
+  f2 = d + c + e
+end
+
+integer(8) function f3 (a1, b1, a2, b2)
+  implicit none
+  integer, value :: a1, b1, a2, b2
+  integer(8) d1, d2
+  !$omp parallel do default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2)
+  do d1 = a1, b1-1
+    do d2 = a2, b2-1
+      u(d1 * 32 + d2) = v(d1 * 32 + d2) + w(d1 * 32 + d2)
+    end do
+  end do
+  f3 = d1 + d2
+end
+end module m
+
+program main
+  use m
+  if (f1 (0, 1024) /= 1024) stop 1
+  if (f2 (0, 1024, 17) /= 1024 + 2 * (17 + 5 * 1024)) stop 2
+  if (f3 (0, 32, 0, 32) /= 64) stop 3
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/pr66199-4.f90 b/libgomp/testsuite/libgomp.fortran/pr66199-4.f90
new file mode 100644
index 0000000..17b62a6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr66199-4.f90
@@ -0,0 +1,60 @@
+! { dg-do run }
+!
+! PR fortran/94690
+! PR middle-end/66199
+
+module m
+  implicit none
+  integer u(0:1023), v(0:1023), w(0:1023)
+  !$omp declare target (u, v, w)
+
+contains
+
+subroutine f1 (a, b)
+  integer a, b, d
+  !$omp target teams distribute parallel do default(none) firstprivate (a, b) shared(u, v, w)
+  do d = a, b-1
+    u(d) = v(d) + w(d)
+  end do
+end
+
+subroutine f2 (a, b, c)
+  integer a, b, c, d, e
+  !$omp target teams distribute parallel do default(none) firstprivate (a, b, c) shared(u, v, w) lastprivate(d, e)
+  do d = a, b-1
+    u(d) = v(d) + w(d)
+    e = c + d * 5
+  end do
+end
+
+subroutine f3 (a1, b1, a2, b2)
+  integer :: a1, b1, a2, b2, d1, d2
+  !$omp target teams distribute parallel do default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) &
+  !$omp&       lastprivate(d1, d2) collapse(2)
+  do d1 = a1, b1-1
+    do d2 = a2, b2-1
+      u(d1 * 32 + d2) = v(d1 * 32 + d2) + w(d1 * 32 + d2)
+    end do
+  end do
+end
+
+subroutine f4 (a1, b1, a2, b2)
+  integer :: a1, b1, a2, b2, d1, d2
+  !$omp target teams distribute parallel do default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) &
+  !$omp&       collapse(2)
+  do d1 = a1, b1-1
+    do d2 = a2, b2-1
+      u(d1 * 32 + d2) = v(d1 * 32 + d2) + w(d1 * 32 + d2)
+    end do
+  end do
+end
+end module m
+
+program main
+  use m
+  implicit none
+  call f1 (0, 1024)
+  call f2 (0, 1024, 17)
+  call f3 (0, 32, 0, 32)
+  call f4 (0, 32, 0, 32)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/pr66199-5.f90 b/libgomp/testsuite/libgomp.fortran/pr66199-5.f90
new file mode 100644
index 0000000..9482f08
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr66199-5.f90
@@ -0,0 +1,71 @@
+! { dg-do run }
+!
+! PR fortran/94690
+! PR middle-end/66199
+
+module m
+  implicit none
+  integer u(0:1023), v(0:1023), w(0:1023)
+  !$omp declare target (u, v, w)
+
+contains
+
+integer function f1 (a, b)
+  integer :: a, b, d
+  !$omp target map(from: d)
+  !$omp teams distribute parallel do simd default(none) firstprivate (a, b) shared(u, v, w)
+  do d = a, b-1
+    u(d) = v(d) + w(d)
+  end do
+  !$omp end target
+  f1 = d
+end
+
+integer function f2 (a, b, c)
+  integer :: a, b, c, d, e
+  !$omp target map(from: d, e)
+  !$omp teams distribute parallel do simd default(none) firstprivate (a, b, c) shared(u, v, w) linear(d) lastprivate(e)
+  do d = a, b-1
+    u(d) = v(d) + w(d)
+    e = c + d * 5
+  end do
+  !$omp end target
+  f2 = d + e
+end
+
+integer function f3 (a1, b1, a2, b2)
+  integer :: a1, b1, a2, b2, d1, d2
+  !$omp target map(from: d1, d2)
+  !$omp teams distribute parallel do simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) &
+  !$omp&      collapse(2)
+  do d1 = a1, b1-1
+    do d2 = a2, b2-1
+      u(d1 * 32 + d2) = v(d1 * 32 + d2) + w(d1 * 32 + d2)
+    end do
+  end do
+  !$omp end target
+  f3 = d1 + d2
+end
+
+integer function f4 (a1, b1, a2, b2)
+  integer :: a1, b1, a2, b2, d1, d2
+  !$omp target map(from: d1, d2)
+  !$omp teams distribute parallel do simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) collapse(2)
+  do d1 = a1, b1-1
+    do d2 = a2, b2-1
+      u(d1 * 32 + d2) = v(d1 * 32 + d2) + w(d1 * 32 + d2)
+    end do
+  end do
+  !$omp end target
+  f4 = d1 + d2
+end
+end module
+
+program main
+  use m
+  implicit none
+  if (f1 (0, 1024) /= 1024) stop 1
+  if (f2 (0, 1024, 17) /= 1024 + (17 + 5 * 1023)) stop 2
+  if (f3 (0, 32, 0, 32) /= 64) stop 3
+  if (f4 (0, 32, 0, 32) /= 64) stop 3
+end
diff --git a/libgomp/testsuite/libgomp.fortran/pr66199-6.f90 b/libgomp/testsuite/libgomp.fortran/pr66199-6.f90
new file mode 100644
index 0000000..f73f683
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr66199-6.f90
@@ -0,0 +1,42 @@
+! { dg-do run }
+!
+! PR fortran/94690
+! PR middle-end/66199
+
+module m
+  implicit none
+  integer :: u(0:1023), v(0:1023), w(0:1023)
+  !$omp declare target (u, v, w)
+
+contains
+
+integer function f2 (a, b, c)
+  integer :: a, b, c, d, e
+  !$omp target map(from: d, e)
+  !$omp teams distribute parallel do default(none) firstprivate (a, b, c) shared(u, v, w) lastprivate(d, e)
+  do d = a, b-1
+    u(d) = v(d) + w(d)
+    e = c + d * 5
+  end do
+  !$omp end target
+  f2 = d + e
+end
+
+integer function f3 (a1, b1, a2, b2)
+  integer :: a1, b1, a2, b2, d1, d2
+  !$omp target map(from: d1, d2)
+  !$omp teams distribute parallel do default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2)
+  do d1 = a1, b1-1
+    do d2 = a2, b2-1
+      u(d1 * 32 + d2) = v(d1 * 32 + d2) + w(d1 * 32 + d2)
+    end do
+  end do
+  !$omp end target
+  f3 = d1 + d2
+end
+end module m
+
+use m
+  if (f2 (0, 1024, 17) /= 1024 + (17 + 5 * 1023)) stop 1
+  if (f3 (0, 32, 0, 32) /= 64) stop 2
+end
diff --git a/libgomp/testsuite/libgomp.fortran/pr66199-7.f90 b/libgomp/testsuite/libgomp.fortran/pr66199-7.f90
new file mode 100644
index 0000000..2bd9468
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr66199-7.f90
@@ -0,0 +1,72 @@
+! { dg-do run }
+!
+! PR fortran/94690
+! PR middle-end/66199
+
+module m
+  implicit none
+  integer u(1024), v(1024), w(1024)
+  !$omp declare target (v, u, w)
+
+contains
+
+integer function f1 (a, b)
+  integer :: a, b, d
+  !$omp target map(from: d)
+  !$omp teams distribute simd default(none) firstprivate (a, b) shared(u, v, w)
+  do d = a, b-1
+    u(d) = v(d) + w(d)
+  end do
+  !$omp end teams distribute simd
+  !$omp end target
+  f1 = d
+end
+
+integer function f2 (a, b, c)
+  integer a, b, c, d, e
+  !$omp target map(from: d, e)
+  !$omp teams distribute simd default(none) firstprivate (a, b, c) shared(u, v, w) linear(d) lastprivate(e)
+  do d = a, b-1
+    u(d) = v(d) + w(d)
+    e = c + d * 5
+  end do
+  !$omp end teams distribute simd
+  !$omp end target
+  f2 = d + e
+end
+
+integer function f3 (a1, b1, a2, b2)
+  integer :: a1, b1, a2, b2, d1, d2
+  !$omp target map(from: d1, d2)
+  !$omp teams distribute simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2)
+  do d1 = a1, b1-1
+    do d2 = a2, b2-1
+      u(d1 * 32 + d2) = v(d1 * 32 + d2) + w(d1 * 32 + d2)
+    end do
+  end do
+  !$omp end teams distribute simd
+  !$omp end target
+  f3 = d1 + d2
+end
+
+integer function f4 (a1, b1, a2, b2)
+  integer :: a1, b1, a2, b2, d1, d2
+  !$omp target map(from: d1, d2)
+  !$omp teams distribute simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) collapse(2)
+  do d1 = a1, b1-1
+    do d2 = a2, b2-1
+      u(d1 * 32 + d2) = v(d1 * 32 + d2) + w(d1 * 32 + d2)
+    end do
+  end do
+  !$omp end teams distribute simd
+  !$omp end target
+  f4 = d1 + d2
+end
+end module
+
+use m
+  if (f1 (0, 1024) /= 1024) stop 1
+  if (f2 (0, 1024, 17) /= 1024 + (17 + 5 * 1023)) stop 2
+  if (f3 (0, 32, 0, 32) /= 64) stop 3
+  if (f4 (0, 32, 0, 32) /= 64) stop 4
+end
diff --git a/libgomp/testsuite/libgomp.fortran/pr66199-8.f90 b/libgomp/testsuite/libgomp.fortran/pr66199-8.f90
new file mode 100644
index 0000000..8a21c6f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr66199-8.f90
@@ -0,0 +1,76 @@
+! { dg-do run }
+!
+! PR fortran/94690
+! PR middle-end/66199
+
+module m
+  implicit none
+  integer u(0:1023), v(0:1023), w(0:1023)
+  !$omp declare target (u, v, w)
+
+contains
+
+integer function f1 (a, b)
+  integer :: a, b, d
+  !$omp target map(from: d)
+  !$omp teams default(none) shared(a, b, d, u, v, w)
+  !$omp distribute simd firstprivate (a, b)
+  do d = a, b-1
+    u(d) = v(d) + w(d)
+  end do
+  !$omp end teams
+  !$omp end target
+  f1 = d
+end
+
+integer function f2 (a, b, c)
+  integer a, b, c, d, e
+  !$omp target map(from: d, e)
+  !$omp teams default(none) firstprivate (a, b, c) shared(d, e, u, v, w)
+  !$omp distribute simd linear(d) lastprivate(e)
+  do d = a, b-1
+    u(d) = v(d) + w(d)
+    e = c + d * 5
+  end do
+  !$omp end teams
+  !$omp end target
+  f2 = d + e
+end
+
+integer function f3 (a1, b1, a2, b2)
+  integer a1, b1, a2, b2, d1, d2
+  !$omp target map(from: d1, d2)
+  !$omp teams default(none) shared(a1, b1, a2, b2, d1, d2, u, v, w)
+  !$omp distribute simd firstprivate (a1, b1, a2, b2) lastprivate(d1, d2) collapse(2)
+  do d1 = a1, b1-1
+    do d2 = a2, b2-1
+      u(d1 * 32 + d2) = v(d1 * 32 + d2) + w(d1 * 32 + d2)
+    end do
+  end do
+  !$omp end teams
+  !$omp end target
+  f3 = d1 + d2
+end
+
+integer function f4 (a1, b1, a2, b2)
+  integer a1, b1, a2, b2, d1, d2
+  !$omp target map(from: d1, d2)
+  !$omp teams default(none) firstprivate (a1, b1, a2, b2) shared(d1, d2, u, v, w)
+  !$omp distribute simd collapse(2)
+  do d1 = a1, b1-1
+    do d2 = a2, b2-1
+      u(d1 * 32 + d2) = v(d1 * 32 + d2) + w(d1 * 32 + d2)
+    end do
+  end do
+  !$omp end teams
+  !$omp end target
+  f4 = d1 + d2
+end
+end module m
+
+use m
+  if (f1 (0, 1024) /= 1024) stop 1
+  if (f2 (0, 1024, 17) /= 1024 + (17 + 5 * 1023)) stop 2
+  if (f3 (0, 32, 0, 32) /= 64) stop 3
+  if (f4 (0, 32, 0, 32) /= 64) stop 4
+end
diff --git a/libgomp/testsuite/libgomp.fortran/pr66199-9.f90 b/libgomp/testsuite/libgomp.fortran/pr66199-9.f90
new file mode 100644
index 0000000..5dde7f8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/pr66199-9.f90
@@ -0,0 +1,46 @@
+! { dg-do run }
+!
+! PR fortran/94690
+! PR middle-end/66199
+
+module m
+  implicit none
+  integer u(1024), v(1024), w(1024)
+  !$omp declare target (u, v, w)
+
+contains
+
+integer function f2 (a, b, c)
+  integer :: a, b, c, d, e
+  !$omp target map(from: d, e)
+  !$omp teams default(none) firstprivate (a, b, c) shared(d, e, u, v, w)
+  !$omp distribute lastprivate(d, e)
+  do d = a, b-1
+    u(d) = v(d) + w(d)
+    e = c + d * 5
+  end do
+  !$omp end teams
+  !$omp end target
+  f2 = d + e
+end
+
+integer function f3 (a1, b1, a2, b2)
+  integer :: a1, b1, a2, b2, d1, d2
+  !$omp target map(from: d1, d2)
+  !$omp teams default(none) shared(a1, b1, a2, b2, d1, d2, u, v, w)
+  !$omp distribute firstprivate (a1, b1, a2, b2) lastprivate(d1, d2) collapse(2)
+  do d1 = a1, b1-1
+    do d2 = a2, b2-1
+      u(d1 * 32 + d2) = v(d1 * 32 + d2) + w(d1 * 32 + d2)
+    end do
+  end do
+  !$omp end teams
+  !$omp end target
+  f3 = d1 + d2
+end
+end module
+
+use m
+  if (f2 (0, 1024, 17) /= 1024 + (17 + 5 * 1023)) stop 1
+  if (f3 (0, 32, 0, 32) /= 64) stop 2
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-1.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-1.f90
new file mode 100644
index 0000000..39faffd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-1.f90
@@ -0,0 +1,38 @@
+! { dg-do run }
+
+program main
+  implicit none
+  integer, allocatable, dimension(:) :: AA, BB, CC, DD
+  integer :: i, N = 20
+
+  allocate(BB(N))
+  AA = [(i, i=1,N)]
+
+  !$omp target enter data map(alloc: BB)
+  !$omp target enter data map(to: AA)
+
+  !$omp target
+    BB = 3 * AA
+  !$omp end target
+
+  !$omp target exit data map(delete: AA)
+  !$omp target exit data map(from: BB)
+
+  if (any (BB /= [(3*i, i=1,N)])) stop 1
+  if (any (AA /= [(i, i=1,N)])) stop 2
+
+
+  CC = 31 * BB
+  DD = [(-i, i=1,N)]
+
+  !$omp target enter data map(to: CC) map(alloc: DD)
+
+  !$omp target
+    DD = 5 * CC
+  !$omp end target
+
+  !$omp target exit data map(delete: CC) map(from: DD)
+
+  if (any (CC /= [(31*3*i, i=1,N)])) stop 3
+  if (any (DD /= [(31*3*5*i, i=1,N)])) stop 4
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-2.F90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-2.F90
new file mode 100644
index 0000000..36a2ed5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-2.F90
@@ -0,0 +1,41 @@
+! { dg-additional-options "-DMEM_SHARED" { target offload_device_shared_as } }
+! { dg-do run }
+!
+! PR middle-end/94635
+  implicit none
+  integer, parameter :: N = 20
+  integer, allocatable, dimension(:) :: my1DPtr
+  integer, dimension(N) :: my1DArr
+  integer :: i
+
+  allocate(my1DPtr(N))
+  my1DPtr = 43
+
+  !$omp target enter data map(alloc: my1DPtr)
+    !$omp target
+      my1DPtr = [(i , i = 1, N)]
+    !$omp end target
+
+    !$omp target map(from: my1DArr) 
+      my1DArr = my1DPtr
+    !$omp end target
+  !$omp target exit data map(delete: my1DPtr)
+
+  if (any (my1DArr /= [(i, i = 1, N)])) stop 1
+#if MEM_SHARED
+  if (any (my1DArr /= my1DPtr)) stop 2
+#else
+  if (any (43 /= my1DPtr)) stop 3
+#endif
+
+  my1DPtr = [(2*N-i, i = 1, N)]
+  my1DArr = 42
+ 
+  !$omp target map(tofrom: my1DArr) map(tofrom: my1DPtr(:))
+    my1DArr = my1DPtr
+    my1DPtr = 20
+  !$omp end target
+
+  if (any (my1DArr /= [(2*N-i, i = 1, N)])) stop 4
+  if (any (20 /= my1DPtr)) stop 6
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target-var.f90 b/libgomp/testsuite/libgomp.fortran/target-var.f90
new file mode 100644
index 0000000..5e5ccd4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-var.f90
@@ -0,0 +1,32 @@
+! { dg-additional-options "-O3" }
+!
+! With -O3 the static local variable A.10 generated for
+! the array constructor [-2, -4, ..., -20] is optimized
+! away - which has to be handled in the offload_vars table.
+!
+program main
+  implicit none (type, external)
+  integer :: j
+  integer, allocatable :: A(:)
+
+  A = [(3*j, j=1, 10)]
+  call bar (A)
+  deallocate (A)
+contains
+  subroutine bar (array)
+    integer :: i
+    integer :: array(:)
+
+    !$omp target map(from:array)
+    !$acc parallel copyout(array)
+    array = [(-2*i, i = 1, size(array))]
+    !$omp do private(array)
+    !$acc loop gang private(array)
+    do i = 1, 10
+      array(i) = 9*i
+    end do
+    if (any (array /= [(-2*i, i = 1, 10)])) error stop 2
+    !$omp end target
+    !$acc end parallel
+  end subroutine bar
+end
diff --git a/libgomp/testsuite/libgomp.fortran/use_device_ptr-optional-2.f90 b/libgomp/testsuite/libgomp.fortran/use_device_ptr-optional-2.f90
index 641ebd9..7a4aaae 100644
--- a/libgomp/testsuite/libgomp.fortran/use_device_ptr-optional-2.f90
+++ b/libgomp/testsuite/libgomp.fortran/use_device_ptr-optional-2.f90
@@ -1,3 +1,4 @@
+! { dg-do run }
 ! Check whether absent optional arguments are properly
 ! handled with use_device_{addr,ptr}.
 program main
diff --git a/libgomp/testsuite/libgomp.oacc-c++/c++.exp b/libgomp/testsuite/libgomp.oacc-c++/c++.exp
index c06c2a0..7200ec1 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/c++.exp
+++ b/libgomp/testsuite/libgomp.oacc-c++/c++.exp
@@ -88,15 +88,6 @@ if { $lang_test_file_found } {
 		unsupported "$subdir $offload_target offloading"
 		continue
 	    }
-	    gcn {
-		if { ![check_effective_target_openacc_amdgcn_accel_present] } {
-		    # Don't bother; execution testing is going to FAIL.
-		    untested "$subdir $offload_target offloading: supported, but hardware not accessible"
-		    continue
-		}
-
-		set acc_mem_shared 0
-	    }
 	    host {
 		set acc_mem_shared 1
 	    }
@@ -115,6 +106,15 @@ if { $lang_test_file_found } {
 
 		set acc_mem_shared 0
 	    }
+	    radeon {
+		if { ![check_effective_target_openacc_radeon_accel_present] } {
+		    # Don't bother; execution testing is going to FAIL.
+		    untested "$subdir $offload_target offloading: supported, but hardware not accessible"
+		    continue
+		}
+
+		set acc_mem_shared 0
+	    }
 	    default {
 		error "Unknown OpenACC device type: $openacc_device_type (offload target: $offload_target)"
 	    }
diff --git a/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
new file mode 100644
index 0000000..ed69359
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
@@ -0,0 +1,58 @@
+#include <openacc.h>
+#include <stdlib.h>
+
+#define N 8
+
+namespace one {
+  int A[N] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+  #pragma acc declare copyin (A)
+};
+
+namespace outer {
+  namespace inner {
+    int B[N];
+    #pragma acc declare create (B)
+  };
+};
+
+static void
+f (void)
+{
+  int i;
+  int C[N];
+  #pragma acc declare copyout (C)
+
+  if (!acc_is_present (&one::A, sizeof (one::A)))
+    abort ();
+
+  if (!acc_is_present (&outer::inner::B, sizeof (outer::inner::B)))
+    abort ();
+
+#pragma acc parallel
+  for (i = 0; i < N; i++)
+    {
+      outer::inner::B[i] = one::A[i];
+      C[i] = outer::inner::B[i];
+    }
+
+#pragma acc parallel
+  for (i = 0; i < N; i++)
+    {
+      if (C[i] != i + 1)
+	abort ();
+    }
+
+#pragma acc parallel
+  for (i = 0; i < N; i++)
+    if (outer::inner::B[i] != i + 1)
+      abort ();
+}
+
+
+int
+main (int argc, char **argv)
+{
+  f ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C b/libgomp/testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C
index c8dba9e..b046bf2 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C
@@ -1,3 +1,12 @@
 /* Verify OpenACC 'firstprivate' mappings for C++ reference types.  */
 
+/* PR middle-end/48591 */
+/* PR other/71064 */
+/* Set to 0 for offloading targets not supporting long double.  */
+#if defined(ACC_DEVICE_TYPE_nvidia) || defined(ACC_DEVICE_TYPE_radeon)
+# define DO_LONG_DOUBLE 0
+#else
+# define DO_LONG_DOUBLE 1
+#endif
+
 #include "../../../gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-3.c
deleted file mode 100644
index 9256500..0000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-3.c
+++ /dev/null
@@ -1,19 +0,0 @@
-/* Test the `acc_get_property' and '`acc_get_property_string' library
-   functions for the host device. */
-/* { dg-additional-sources acc_get_property-aux.c } */
-/* { dg-do run } */
-
-#include <openacc.h>
-#include <stdio.h>
-
-void expect_device_properties
-(acc_device_t dev_type, int dev_num,
- int expected_total_mem, int expected_free_mem,
- const char* expected_vendor, const char* expected_name,
- const char* expected_driver);
-
-int main()
-{
-  printf ("Checking acc_device_host device properties\n");
-  expect_device_properties (acc_device_host, 0, 0, 0, "GNU", "GOMP", "1.0");
-}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-aux.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-aux.c
index 952bdbf..47285fc 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-aux.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-aux.c
@@ -6,11 +6,12 @@
 #include <stdio.h>
 #include <string.h>
 
-void expect_device_properties
-(acc_device_t dev_type, int dev_num,
- int expected_total_mem, int expected_free_mem,
- const char* expected_vendor, const char* expected_name,
- const char* expected_driver)
+
+void
+expect_device_string_properties (acc_device_t dev_type, int dev_num,
+				 const char* expected_vendor,
+				 const char* expected_name,
+				 const char* expected_driver)
 {
   const char *vendor = acc_get_property_string (dev_num, dev_type,
 						acc_property_vendor);
@@ -21,25 +22,6 @@ void expect_device_properties
       abort ();
     }
 
-  int total_mem = acc_get_property (dev_num, dev_type,
-				    acc_property_memory);
-  if (total_mem != expected_total_mem)
-    {
-      fprintf (stderr, "Expected acc_property_memory to equal %d, "
-	       "but was %d.\n", expected_total_mem, total_mem);
-      abort ();
-
-    }
-
-  int free_mem = acc_get_property (dev_num, dev_type,
-				   acc_property_free_memory);
-  if (free_mem != expected_free_mem)
-    {
-      fprintf (stderr, "Expected acc_property_free_memory to equal %d, "
-	       "but was %d.\n", expected_free_mem, free_mem);
-      abort ();
-    }
-
   const char *name = acc_get_property_string (dev_num, dev_type,
 					      acc_property_name);
   if (strcmp (name, expected_name))
@@ -59,11 +41,11 @@ void expect_device_properties
     }
 
   int unknown_property = 16058;
-  int v = acc_get_property (dev_num, dev_type, (acc_device_property_t)unknown_property);
+  size_t v = acc_get_property (dev_num, dev_type, (acc_device_property_t)unknown_property);
   if (v != 0)
     {
       fprintf (stderr, "Expected value of unknown numeric property to equal 0, "
-	       "but was %d.\n", v);
+	       "but was %zu.\n", v);
       abort ();
     }
 
@@ -72,9 +54,45 @@ void expect_device_properties
   if (s != NULL)
     {
       fprintf (stderr, "Expected value of unknown string property to be NULL, "
-	       "but was %d.\n", s);
+	       "but was %s.\n", s);
       abort ();
     }
+}
 
+void
+expect_device_memory (acc_device_t dev_type, int dev_num,
+		      size_t expected_total_memory)
+{
 
+  size_t total_mem = acc_get_property (dev_num, dev_type,
+				       acc_property_memory);
+
+  if (total_mem != expected_total_memory)
+    {
+      fprintf (stderr, "Expected acc_property_memory to equal %zu, "
+	       "but was %zu.\n", expected_total_memory, total_mem);
+      abort ();
+    }
+
+  size_t free_mem = acc_get_property (dev_num, dev_type,
+				      acc_property_free_memory);
+  if (free_mem > total_mem)
+    {
+      fprintf (stderr, "Expected acc_property_free_memory <= acc_property_memory"
+	       ", but free memory was %zu and total memory was %zu.\n",
+	       free_mem, total_mem);
+      abort ();
+    }
+}
+
+void
+expect_device_properties (acc_device_t dev_type, int dev_num,
+			  size_t expected_total_memory,
+			  const char* expected_vendor,
+			  const char* expected_name,
+			  const char* expected_driver)
+{
+  expect_device_string_properties (dev_type, dev_num, expected_vendor,
+				   expected_name, expected_driver);
+  expect_device_memory (dev_type, dev_num, expected_total_memory);
 }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-gcn.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-gcn.c
new file mode 100644
index 0000000..4b1fb5e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-gcn.c
@@ -0,0 +1,135 @@
+/* Test the `acc_get_property' and `acc_get_property_string' library
+   functions on amdgcn devices by comparing property values with
+   those obtained through the HSA API. */
+/* { dg-additional-sources acc_get_property-aux.c } */
+/* { dg-additional-options "-ldl" } */
+/* { dg-do run { target openacc_radeon_accel_selected } } */
+
+#include <dlfcn.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <openacc.h>
+
+#ifndef __cplusplus
+typedef int bool;
+#endif
+#include <hsa.h>
+
+
+void expect_device_string_properties (acc_device_t dev_type, int dev_num,
+				      const char* expected_vendor,
+				      const char* expected_name,
+				      const char* expected_driver);
+
+hsa_status_t (*hsa_agent_get_info_fn) (hsa_agent_t agent,
+				       hsa_agent_info_t attribute,
+				       void *value);
+hsa_status_t (*hsa_system_get_info_fn) (hsa_system_info_t attribute,
+					void *value);
+hsa_status_t (*hsa_iterate_agents_fn)
+(hsa_status_t (*callback)(hsa_agent_t agent, void *data), void *data);
+hsa_status_t (*hsa_init_fn) (void);
+
+char* support_cpu_devices;
+
+void
+test_setup ()
+{
+  char* env_runtime;
+  char* hsa_runtime_lib;
+  void *handle;
+
+#define DLSYM_FN(function)						\
+  function##_fn = (typeof(function##_fn))dlsym (handle, #function);	\
+  if (function##_fn == NULL)						\
+    {									\
+      fprintf (stderr, "Could not get symbol " #function ".\n");	\
+      abort (); 							\
+    }
+
+  env_runtime = getenv ("HSA_RUNTIME_LIB");
+  hsa_runtime_lib = env_runtime ? env_runtime : (char*)"libhsa-runtime64.so";
+
+  handle = dlopen (hsa_runtime_lib, RTLD_LAZY);
+  if (!handle)
+    {
+      fprintf (stderr, "Could not load %s.\n", hsa_runtime_lib);
+      abort ();
+    }
+
+  DLSYM_FN (hsa_agent_get_info)
+  DLSYM_FN (hsa_system_get_info)
+  DLSYM_FN (hsa_iterate_agents)
+  DLSYM_FN (hsa_init)
+
+  hsa_init_fn ();
+
+  support_cpu_devices = getenv ("GCN_SUPPORT_CPU_DEVICES");
+}
+
+static hsa_status_t
+check_agent_properties (hsa_agent_t agent, void *dev_num_arg)
+{
+
+  char name[64];
+  char vendor_name[64];
+  uint16_t minor;
+  uint16_t major;
+  char driver[60];
+
+  hsa_status_t status;
+  hsa_device_type_t device_type;
+  int* dev_num = (int*)dev_num_arg;
+
+#define AGENT_GET_INFO(info_type, val)				\
+  status = hsa_agent_get_info_fn (agent, info_type, &val);	\
+  if (status != HSA_STATUS_SUCCESS)				\
+    {								\
+      fprintf (stderr, "Failed to obtain " #info_type ".\n");	\
+      abort ();							\
+    }
+#define SYSTEM_GET_INFO(info_type, val)				\
+  status = hsa_system_get_info_fn (info_type, &val);		\
+  if (status != HSA_STATUS_SUCCESS)				\
+    {								\
+      fprintf (stderr, "Failed to obtain " #info_type ".\n");	\
+      abort ();							\
+    }
+
+  AGENT_GET_INFO (HSA_AGENT_INFO_DEVICE, device_type)
+
+    /* Skip unsupported device types.  Mimic the GCN plugin's behavior. */
+    if (!(device_type == HSA_DEVICE_TYPE_GPU
+	  || (support_cpu_devices && device_type == HSA_DEVICE_TYPE_CPU)))
+      return HSA_STATUS_SUCCESS;
+
+  AGENT_GET_INFO (HSA_AGENT_INFO_NAME, name)
+  AGENT_GET_INFO (HSA_AGENT_INFO_VENDOR_NAME, vendor_name)
+
+  SYSTEM_GET_INFO (HSA_SYSTEM_INFO_VERSION_MINOR, minor)
+  SYSTEM_GET_INFO (HSA_SYSTEM_INFO_VERSION_MAJOR, major)
+
+  snprintf (driver, sizeof driver, "HSA Runtime %hu.%hu",
+	    (unsigned short int)major, (unsigned short int)minor);
+
+  expect_device_string_properties(acc_device_radeon, *dev_num,
+				  vendor_name, name, driver);
+
+  (*dev_num)++;
+
+  return status;
+}
+
+int
+main ()
+{
+  int dev_num = 0;
+  test_setup ();
+
+  hsa_status_t status =
+    hsa_iterate_agents_fn (&check_agent_properties, &dev_num);
+
+  return status;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-host.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-host.c
new file mode 100644
index 0000000..4ed0dfa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-host.c
@@ -0,0 +1,20 @@
+/* Test the `acc_get_property' and '`acc_get_property_string' library
+   functions for the host device. */
+/* { dg-additional-sources acc_get_property-aux.c } */
+/* { dg-do run } */
+
+#include <openacc.h>
+#include <stdio.h>
+
+void expect_device_properties (acc_device_t dev_type, int dev_num,
+			       size_t expected_memory,
+			       const char* expected_vendor,
+			       const char* expected_name,
+			       const char* expected_driver);
+
+int
+main ()
+{
+  printf ("Checking acc_device_host device properties\n");
+  expect_device_properties (acc_device_host, 0, 0, "GNU", "GOMP", "1.0");
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-nvptx.c
index 4dd13c4..6334cfd 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property-nvptx.c
@@ -11,13 +11,14 @@
 #include <string.h>
 #include <stdio.h>
 
-void expect_device_properties
-(acc_device_t dev_type, int dev_num,
- int expected_total_mem, int expected_free_mem,
- const char* expected_vendor, const char* expected_name,
- const char* expected_driver);
+void expect_device_properties (acc_device_t dev_type, int dev_num,
+			       size_t expected_memory,
+			       const char* expected_vendor,
+			       const char* expected_name,
+			       const char* expected_driver);
 
-int main ()
+int
+main ()
 {
   int dev_count;
   cudaGetDeviceCount (&dev_count);
@@ -30,26 +31,26 @@ int main ()
 	  abort ();
 	}
 
-      printf("Checking device %d\n", dev_num);
+      printf ("Checking device %d\n", dev_num);
 
       const char *vendor = "Nvidia";
       size_t free_mem;
       size_t total_mem;
-      if (cudaMemGetInfo(&free_mem, &total_mem) != cudaSuccess)
+      if (cudaMemGetInfo (&free_mem, &total_mem) != cudaSuccess)
 	{
 	  fprintf (stderr, "cudaMemGetInfo failed.\n");
 	  abort ();
 	}
 
       struct cudaDeviceProp p;
-      if (cudaGetDeviceProperties(&p, dev_num) != cudaSuccess)
+      if (cudaGetDeviceProperties (&p, dev_num) != cudaSuccess)
 	{
 	  fprintf (stderr, "cudaGetDeviceProperties failed.\n");
 	  abort ();
 	}
 
       int driver_version;
-      if (cudaDriverGetVersion(&driver_version) != cudaSuccess)
+      if (cudaDriverGetVersion (&driver_version) != cudaSuccess)
 	{
 	  fprintf (stderr, "cudaDriverGetVersion failed.\n");
 	  abort ();
@@ -62,7 +63,9 @@ int main ()
       snprintf (driver, sizeof driver, "CUDA Driver %u.%u",
 		driver_version / 1000, driver_version % 1000 / 10);
 
-      expect_device_properties(acc_device_nvidia, dev_num,
-			       total_mem, free_mem, vendor, p.name, driver);
+      /* Note that this check relies on the fact that the device numbering
+	 used by the nvptx plugin agrees with the CUDA device ordering. */
+      expect_device_properties (acc_device_nvidia, dev_num,
+				total_mem, vendor, p.name, driver);
     }
 }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property.c
index 289d1ba..3460035 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_get_property.c
@@ -3,8 +3,7 @@
    of all device types mentioned in the OpenACC standard.
 
    See also acc_get_property.f90. */
-/* { dg-do run { target { { ! { openacc_host_selected } } && { ! { openacc_amdgcn_accel_selected } } } } } */
-/* FIXME: This test does not work with the GCN implementation stub yet.  */
+/* { dg-do run } */
 
 #include <openacc.h>
 #include <stdlib.h>
@@ -15,16 +14,16 @@
    and do basic device independent validation. */
 
 void
-print_device_properties(acc_device_t type)
+print_device_properties (acc_device_t type)
 {
   const char *s;
   size_t v;
 
-  int dev_count = acc_get_num_devices(type);
+  int dev_count = acc_get_num_devices (type);
 
   for (int i = 0; i < dev_count; ++i)
     {
-      printf("  Device %d:\n", i+1);
+      printf ("  Device %d:\n", i+1);
 
       s = acc_get_property_string (i, type, acc_property_vendor);
       printf ("    Vendor: %s\n", s);
@@ -35,10 +34,10 @@ print_device_properties(acc_device_t type)
 	}
 
       v = acc_get_property (i, type,  acc_property_memory);
-      printf ("    Total memory: %zd\n", v);
+      printf ("    Total memory: %zu\n", v);
 
       v = acc_get_property (i, type, acc_property_free_memory);
-      printf ("    Free memory: %zd\n", v);
+      printf ("    Free memory: %zu\n", v);
 
       s = acc_get_property_string (i, type, acc_property_name);
       printf ("    Name: %s\n", s);
@@ -58,19 +57,20 @@ print_device_properties(acc_device_t type)
     }
 }
 
-int main ()
+int
+main ()
 {
-  printf("acc_device_none:\n");
+  printf ("acc_device_none:\n");
   /* For completness; not expected to print anything since there
      should be no devices of this type. */
-  print_device_properties(acc_device_none);
+  print_device_properties (acc_device_none);
 
-  printf("acc_device_default:\n");
-  print_device_properties(acc_device_default);
+  printf ("acc_device_default:\n");
+  print_device_properties (acc_device_default);
 
-  printf("acc_device_host:\n");
-  print_device_properties(acc_device_host);
+  printf ("acc_device_host:\n");
+  print_device_properties (acc_device_host);
 
-  printf("acc_device_not_host:\n");
-  print_device_properties(acc_device_not_host);
+  printf ("acc_device_not_host:\n");
+  print_device_properties (acc_device_not_host);
 }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c
index e82a03e..7d05f48 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c
@@ -224,7 +224,7 @@ static void cb_compute_construct_end (acc_prof_info *prof_info, acc_event_info *
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
-  else if (acc_device_type == acc_device_gcn)
+  else if (acc_device_type == acc_device_radeon)
     assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c
index ddf647c..ad33f72 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c
@@ -106,7 +106,7 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e
     assert (event_info->launch_event.vector_length >= 1);
   else if (acc_device_type == acc_device_nvidia) /* ... is special.  */
     assert (event_info->launch_event.vector_length == 32);
-  else if (acc_device_type == acc_device_gcn) /* ...and so is this.  */
+  else if (acc_device_type == acc_device_radeon) /* ...and so is this.  */
     assert (event_info->launch_event.vector_length == 64);
   else
     {
@@ -120,7 +120,7 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
-  else if (acc_device_type == acc_device_gcn)
+  else if (acc_device_type == acc_device_radeon)
     assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c
index dc7c758..a5e9ab3 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c
@@ -265,7 +265,7 @@ static void cb_enter_data_end (acc_prof_info *prof_info, acc_event_info *event_i
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
-  else if (acc_device_type == acc_device_gcn)
+  else if (acc_device_type == acc_device_radeon)
     assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
@@ -321,7 +321,7 @@ static void cb_exit_data_start (acc_prof_info *prof_info, acc_event_info *event_
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
-  else if (acc_device_type == acc_device_gcn)
+  else if (acc_device_type == acc_device_radeon)
     assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
@@ -375,7 +375,7 @@ static void cb_exit_data_end (acc_prof_info *prof_info, acc_event_info *event_in
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
-  else if (acc_device_type == acc_device_gcn)
+  else if (acc_device_type == acc_device_radeon)
     assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
@@ -516,7 +516,7 @@ static void cb_compute_construct_end (acc_prof_info *prof_info, acc_event_info *
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
-  else if (acc_device_type == acc_device_gcn)
+  else if (acc_device_type == acc_device_radeon)
     assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
@@ -581,7 +581,7 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *e
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
-  else if (acc_device_type == acc_device_gcn)
+  else if (acc_device_type == acc_device_radeon)
     assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
@@ -647,7 +647,7 @@ static void cb_enqueue_launch_end (acc_prof_info *prof_info, acc_event_info *eve
 
   if (acc_device_type == acc_device_host)
     assert (api_info->device_api == acc_device_api_none);
-  else if (acc_device_type == acc_device_gcn)
+  else if (acc_device_type == acc_device_radeon)
     assert (api_info->device_api == acc_device_api_other);
   else
     assert (api_info->device_api == acc_device_api_cuda);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c
index 840052f..7496426 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c
@@ -26,8 +26,8 @@ main ()
   acc_device_t d;
 #if defined ACC_DEVICE_TYPE_nvidia
   d = acc_device_nvidia;
-#elif defined ACC_DEVICE_TYPE_gcn
-  d = acc_device_gcn;
+#elif defined ACC_DEVICE_TYPE_radeon
+  d = acc_device_radeon;
 #elif defined ACC_DEVICE_TYPE_host
   d = acc_device_host;
 #else
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c
index a59047a..13e5ca2 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c
@@ -38,7 +38,7 @@ main ()
 	assert (v.b[i] == v.a + i);
 
       assert (!acc_is_present (&v, sizeof (v)));
-      assert (!acc_is_present (v.b, sizeof (int *) * n));
+      assert (!acc_is_present (v.b, sizeof (int) * n));
     }
 
   return 0;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c
index 0ca5990..1b4cf2f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c
@@ -41,9 +41,9 @@ main ()
 	assert (v.b[i] == v.a + i);
 
       assert (acc_is_present (&v, sizeof (v)));
-      assert (!acc_is_present (v.b, sizeof (int *) * n));
-      assert (!acc_is_present (v.c, sizeof (int *) * n));
-      assert (!acc_is_present (v.d, sizeof (int *) * n));
+      assert (!acc_is_present (v.b, sizeof (int) * n));
+      assert (!acc_is_present (v.c, sizeof (int) * n));
+      assert (!acc_is_present (v.d, sizeof (int) * n));
     }
 
 #pragma acc exit data copyout(v)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c
index 4a8b310..2cdd2d1 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c
@@ -3,4 +3,13 @@
 /* { dg-additional-options "-Wno-psabi" } as apparently we're doing funny
    things with vector arguments.  */
 
+/* PR middle-end/48591 */
+/* PR other/71064 */
+/* Set to 0 for offloading targets not supporting long double.  */
+#if defined(ACC_DEVICE_TYPE_nvidia) || defined(ACC_DEVICE_TYPE_radeon)
+# define DO_LONG_DOUBLE 0
+#else
+# define DO_LONG_DOUBLE 1
+#endif
+
 #include "../../../gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c
index 517004a..64f8ab8 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c
@@ -1,11 +1,11 @@
 /* { dg-do link } */
-/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" { target { openacc_nvidia_accel_selected || openacc_amdgcn_accel_selected } } } */
+/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" { target { openacc_nvidia_accel_selected || openacc_radeon_accel_selected } } } */
 
 int var;
 #pragma acc declare create (var)
 
 void __attribute__((noinline, noclone))
-foo () /* { dg-error "function 'foo' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" "" { target { openacc_nvidia_accel_selected || openacc_amdgcn_accel_selected } } } */
+foo () /* { dg-error "function 'foo' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" "" { target { openacc_nvidia_accel_selected || openacc_radeon_accel_selected } } } */
 {
   var++;
 }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-7.c
new file mode 100644
index 0000000..6830ef1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-7.c
@@ -0,0 +1,66 @@
+/* { dg-do run } */
+
+/* Test if, if_present clauses on host_data construct.  */
+/* C/C++ variant of 'libgomp.oacc-fortran/host_data-5.F90' */
+
+#include <assert.h>
+#include <stdint.h>
+
+void
+foo (float *p, intptr_t host_p, int cond)
+{
+  assert (p == (float *) host_p);
+
+#pragma acc data copyin(host_p)
+  {
+#pragma acc host_data use_device(p) if_present
+    /* p not mapped yet, so it will be equal to the host pointer.  */
+    assert (p == (float *) host_p);
+
+#pragma acc data copy(p[0:100])
+    {
+      /* Not inside a host_data construct, so p is still the host pointer.  */
+      assert (p == (float *) host_p);
+
+#pragma acc host_data use_device(p)
+      {
+#if ACC_MEM_SHARED
+	assert (p == (float *) host_p);
+#else
+	/* The device address is different from host address.  */
+	assert (p != (float *) host_p);
+#endif
+      }
+
+#pragma acc host_data use_device(p) if_present
+      {
+#if ACC_MEM_SHARED
+	assert (p == (float *) host_p);
+#else
+	/* p is present now, so this is the same as above.  */
+	assert (p != (float *) host_p);
+#endif
+      }
+
+#pragma acc host_data use_device(p) if(cond)
+      {
+#if ACC_MEM_SHARED
+	assert (p == (float *) host_p);
+#else
+	/* p is the device pointer iff cond is true.  */
+	assert ((p != (float *) host_p) == cond);
+#endif
+      }
+    }
+  }
+}
+
+int
+main (void)
+{
+  float arr[100];
+  foo (arr, (intptr_t) arr, 0);
+  foo (arr, (intptr_t) arr, 1);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
index 34bc57e..0273c2b 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
@@ -1,3 +1,6 @@
+/* AMD GCN does not use 32-lane vectors.
+   { dg-skip-if "unsuitable dimensions" { openacc_radeon_accel_selected } { "*" } { "" } } */
+
 /* { dg-additional-options "-fopenacc-dim=32" } */
 
 #include <stdio.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c
index 04387d3..ca77164 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c
@@ -128,5 +128,14 @@ int test_1 (int gp, int wp, int vp)
 
 int main ()
 {
+#ifdef ACC_DEVICE_TYPE_radeon
+  /* AMD GCN uses the autovectorizer for the vector dimension: the use
+     of a function call in vector-partitioned code in this test is not
+     currently supported.  */
+  /* AMD GCN does not currently support multiple workers.  This should be
+     set to 16 when that changes.  */
+  return test_1 (16, 1, 1);
+#else
   return test_1 (16, 16, 32);
+#endif
 }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c
index 766e578..5c84301 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c
@@ -9,11 +9,13 @@ int main ()
   int ix;
   int exit = 0;
   int ondev = 0;
+  int gangsize, workersize, vectorsize;
 
   for (ix = 0; ix < N;ix++)
     ary[ix] = -1;
   
-#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) copy(ary) copy(ondev)
+#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
+	    copy(ary) copy(ondev) copyout(gangsize, workersize, vectorsize)
   {
 #pragma acc loop gang worker vector
     for (unsigned ix = 0; ix < N; ix++)
@@ -32,6 +34,10 @@ int main ()
 	else
 	  ary[ix] = ix;
       }
+
+    gangsize = __builtin_goacc_parlevel_size (GOMP_DIM_GANG);
+    workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER);
+    vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR);
   }
 
   for (ix = 0; ix < N; ix++)
@@ -39,11 +45,12 @@ int main ()
       int expected = ix;
       if(ondev)
 	{
-	  int chunk_size = (N + 32*32*32 - 1) / (32*32*32);
+	  int chunk_size = (N + gangsize * workersize * vectorsize - 1)
+			   / (gangsize * workersize * vectorsize);
 	  
-	  int g = ix / (chunk_size * 32 * 32);
-	  int w = ix / 32 % 32;
-	  int v = ix % 32;
+	  int g = ix / (chunk_size * workersize * vectorsize);
+	  int w = (ix / vectorsize) % workersize;
+	  int v = ix % vectorsize;
 
 	  expected = (g << 16) | (w << 8) | v;
 	}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c
index 0bec6e1..9c4a85f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c
@@ -8,8 +8,10 @@ int main ()
   int ix;
   int ondev = 0;
   int t = 0, h = 0;
-  
-#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) copy(ondev)
+  int gangsize, workersize, vectorsize;
+
+#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
+	copy(ondev) copyout(gangsize, workersize, vectorsize)
   {
 #pragma acc loop gang worker vector reduction(+:t)
     for (unsigned ix = 0; ix < N; ix++)
@@ -28,18 +30,22 @@ int main ()
 	  }
 	t += val;
       }
+    gangsize = __builtin_goacc_parlevel_size (GOMP_DIM_GANG);
+    workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER);
+    vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR);
   }
 
   for (ix = 0; ix < N; ix++)
     {
       int val = ix;
-      if(ondev)
+      if (ondev)
 	{
-	  int chunk_size = (N + 32*32*32 - 1) / (32*32*32);
+	  int chunk_size = (N + gangsize * workersize * vectorsize - 1)
+			   / (gangsize * workersize * vectorsize);
 	  
-	  int g = ix / (chunk_size * 32 * 32);
-	  int w = ix / 32 % 32;
-	  int v = ix % 32;
+	  int g = ix / (chunk_size * vectorsize * workersize);
+	  int w = ix / vectorsize % workersize;
+	  int v = ix % vectorsize;
 
 	  val = (g << 16) | (w << 8) | v;
 	}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c
index da4921d..1173c1f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c
@@ -9,8 +9,9 @@ int main ()
   int ix;
   int ondev = 0;
   int t = 0,  h = 0;
+  int vectorsize;
 
-#pragma acc parallel vector_length(32) copy(ondev)
+#pragma acc parallel vector_length(32) copy(ondev) copyout(vectorsize)
   {
 #pragma acc loop vector reduction (+:t)
     for (unsigned ix = 0; ix < N; ix++)
@@ -29,6 +30,7 @@ int main ()
 	  }
 	t += val;
       }
+    vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR);
   }
 
   for (ix = 0; ix < N; ix++)
@@ -38,7 +40,7 @@ int main ()
 	{
 	  int g = 0;
 	  int w = 0;
-	  int v = ix % 32;
+	  int v = ix % vectorsize;
 
 	  val = (g << 16) | (w << 8) | v;
 	}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c
index 15e2bc2..84c2296 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c
@@ -9,8 +9,9 @@ int main ()
   int ix;
   int ondev = 0;
   int q = 0,  h = 0;
+  int vectorsize;
 
-#pragma acc parallel vector_length(32) copy(q) copy(ondev)
+#pragma acc parallel vector_length(32) copy(q) copy(ondev) copyout(vectorsize)
   {
     int t = q;
     
@@ -32,6 +33,7 @@ int main ()
 	t += val;
       }
     q = t;
+    vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR);
   }
 
   for (ix = 0; ix < N; ix++)
@@ -41,7 +43,7 @@ int main ()
 	{
 	  int g = 0;
 	  int w = 0;
-	  int v = ix % 32;
+	  int v = ix % vectorsize;
 
 	  val = (g << 16) | (w << 8) | v;
 	}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
index 6bbd04f..648f89e 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
@@ -8,8 +8,10 @@ int main ()
   int ix;
   int ondev = 0;
   int t = 0,  h = 0;
+  int workersize;
 
-#pragma acc parallel num_workers(32) vector_length(32) copy(ondev)
+#pragma acc parallel num_workers(32) vector_length(32) copy(ondev) \
+	    copyout(workersize)
   {
 #pragma acc loop worker reduction(+:t)
     for (unsigned ix = 0; ix < N; ix++)
@@ -28,6 +30,7 @@ int main ()
 	  }
 	t += val;
       }
+    workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER);
   }
 
   for (ix = 0; ix < N; ix++)
@@ -36,7 +39,7 @@ int main ()
       if(ondev)
 	{
 	  int g = 0;
-	  int w = ix % 32;
+	  int w = ix % workersize;
 	  int v = 0;
 
 	  val = (g << 16) | (w << 8) | v;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
index c63a5d4..f9fcf37 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
@@ -8,8 +8,10 @@ int main ()
   int ix;
   int ondev = 0;
   int q = 0,  h = 0;
+  int workersize;
 
-#pragma acc parallel num_workers(32) vector_length(32) copy(q) copy(ondev)
+#pragma acc parallel num_workers(32) vector_length(32) copy(q) copy(ondev) \
+	    copyout(workersize)
   {
     int t = q;
     
@@ -31,6 +33,7 @@ int main ()
 	t += val;
       }
     q = t;
+    workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER);
   }
 
   for (ix = 0; ix < N; ix++)
@@ -39,7 +42,7 @@ int main ()
       if(ondev)
 	{
 	  int g = 0;
-	  int w = ix % 32;
+	  int w = ix % workersize;
 	  int v = 0;
 
 	  val = (g << 16) | (w << 8) | v;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c
index 71d3969..c360ad1 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c
@@ -8,8 +8,10 @@ int main ()
   int ix;
   int ondev = 0;
   int t = 0, h = 0;
+  int workersize, vectorsize;
   
-#pragma acc parallel num_workers(32) vector_length(32) copy(ondev)
+#pragma acc parallel num_workers(32) vector_length(32) copy(ondev) \
+	    copyout(workersize, vectorsize)
   {
 #pragma acc loop worker vector reduction (+:t)
     for (unsigned ix = 0; ix < N; ix++)
@@ -28,6 +30,8 @@ int main ()
 	  }
 	t += val;
       }
+    workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER);
+    vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR);
   }
 
   for (ix = 0; ix < N; ix++)
@@ -36,8 +40,8 @@ int main ()
       if(ondev)
 	{
 	  int g = 0;
-	  int w = (ix / 32) % 32;
-	  int v = ix % 32;
+	  int w = (ix / vectorsize) % workersize;
+	  int v = ix % vectorsize;
 
 	  val = (g << 16) | (w << 8) | v;
 	}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c
index 6010cd2..8c858f3 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c
@@ -9,11 +9,13 @@ int main ()
   int ix;
   int exit = 0;
   int ondev = 0;
+  int vectorsize;
 
   for (ix = 0; ix < N;ix++)
     ary[ix] = -1;
   
-#pragma acc parallel vector_length(32) copy(ary) copy(ondev)
+#pragma acc parallel vector_length(32) copy(ary) copy(ondev) \
+	    copyout(vectorsize)
   {
 #pragma acc loop vector
     for (unsigned ix = 0; ix < N; ix++)
@@ -31,6 +33,7 @@ int main ()
 	else
 	  ary[ix] = ix;
       }
+    vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR);
   }
 
   for (ix = 0; ix < N; ix++)
@@ -40,7 +43,7 @@ int main ()
 	{
 	  int g = 0;
 	  int w = 0;
-	  int v = ix % 32;
+	  int v = ix % vectorsize;
 
 	  expected = (g << 16) | (w << 8) | v;
 	}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
index fa6fb91..5fe486f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
@@ -9,11 +9,13 @@ int main ()
   int ix;
   int exit = 0;
   int ondev = 0;
+  int workersize;
 
   for (ix = 0; ix < N;ix++)
     ary[ix] = -1;
   
-#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev)
+#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) \
+	    copyout(workersize)
   {
 #pragma acc loop worker
     for (unsigned ix = 0; ix < N; ix++)
@@ -31,6 +33,7 @@ int main ()
 	else
 	  ary[ix] = ix;
       }
+    workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER);
   }
 
   for (ix = 0; ix < N; ix++)
@@ -39,7 +42,7 @@ int main ()
       if(ondev)
 	{
 	  int g = 0;
-	  int w = ix % 32;
+	  int w = ix % workersize;
 	  int v = 0;
 
 	  expected = (g << 16) | (w << 8) | v;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c
index cd4cc99..fd4e4cf 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c
@@ -9,11 +9,13 @@ int main ()
   int ix;
   int exit = 0;
   int ondev = 0;
+  int workersize, vectorsize;
 
   for (ix = 0; ix < N;ix++)
     ary[ix] = -1;
   
-#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev)
+#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) \
+	    copyout(workersize, vectorsize)
   {
 #pragma acc loop worker vector
     for (unsigned ix = 0; ix < N; ix++)
@@ -31,6 +33,8 @@ int main ()
 	else
 	  ary[ix] = ix;
       }
+    workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER);
+    vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR);
   }
 
   for (ix = 0; ix < N; ix++)
@@ -39,8 +43,8 @@ int main ()
       if(ondev)
 	{
 	  int g = 0;
-	  int w = (ix / 32) % 32;
-	  int v = ix % 32;
+	  int w = (ix / vectorsize) % workersize;
+	  int v = ix % vectorsize;
 
 	  expected = (g << 16) | (w << 8) | v;
 	}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
index a5edfc6..cc4c738 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
@@ -14,7 +14,8 @@ static unsigned int __attribute__ ((optimize ("O2"))) acc_gang ()
 {
   if (acc_on_device ((int) acc_device_host))
     return 0;
-  else if (acc_on_device ((int) acc_device_nvidia))
+  else if (acc_on_device ((int) acc_device_nvidia)
+	   || acc_on_device ((int) acc_device_radeon))
     return __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
   else
     __builtin_abort ();
@@ -25,7 +26,8 @@ static unsigned int __attribute__ ((optimize ("O2"))) acc_worker ()
 {
   if (acc_on_device ((int) acc_device_host))
     return 0;
-  else if (acc_on_device ((int) acc_device_nvidia))
+  else if (acc_on_device ((int) acc_device_nvidia)
+	   || acc_on_device ((int) acc_device_radeon))
     return __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
   else
     __builtin_abort ();
@@ -36,7 +38,8 @@ static unsigned int __attribute__ ((optimize ("O2"))) acc_vector ()
 {
   if (acc_on_device ((int) acc_device_host))
     return 0;
-  else if (acc_on_device ((int) acc_device_nvidia))
+  else if (acc_on_device ((int) acc_device_nvidia)
+	   || acc_on_device ((int) acc_device_radeon))
     return __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
   else
     __builtin_abort ();
@@ -282,6 +285,12 @@ int main ()
 	  /* The GCC nvptx back end enforces num_workers (32).  */
 	  workers_actual = 32;
 	}
+      else if (acc_on_device (acc_device_radeon))
+	{
+	  /* The GCC GCN back end is limited to num_workers (16).
+	     Temporarily set this to 1 until multiple workers are permitted. */
+	  workers_actual = 1; // 16;
+	}
       else
 	__builtin_abort ();
 #pragma acc loop worker reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
@@ -328,6 +337,11 @@ int main ()
 	  /* We're actually executing with num_workers (32).  */
 	  /* workers_actual = 32; */
 	}
+      else if (acc_on_device (acc_device_radeon))
+	{
+	  /* The GCC GCN back end is limited to num_workers (16).  */
+	  workers_actual = 16;
+	}
       else
 	__builtin_abort ();
 #pragma acc loop worker reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
@@ -367,6 +381,11 @@ int main ()
 	  /* The GCC nvptx back end enforces vector_length (32).  */
 	  vectors_actual = 1024;
 	}
+      else if (acc_on_device (acc_device_radeon))
+	{
+	  /* The GCC GCN back end enforces vector_length (1): autovectorize. */
+	  vectors_actual = 1;
+	}
       else
 	__builtin_abort ();
 #pragma acc loop vector reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
@@ -407,6 +426,13 @@ int main ()
 	  /* The GCC nvptx back end enforces vector_length (32).  */
 	  vectors_actual = 32;
 	}
+      else if (acc_on_device (acc_device_radeon))
+	{
+	  /* Because of the way vectors are implemented for GCN, a vector loop
+	     containing a seq routine call will not vectorize calls to that
+	     routine.  Hence, we'll only get one "vector".  */
+	  vectors_actual = 1;
+	}
       else
 	__builtin_abort ();
 #pragma acc loop vector reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
@@ -433,6 +459,9 @@ int main ()
        in the following case.  So, limit ourselves here.  */
     if (acc_get_device_type () == acc_device_nvidia)
       gangs = 3;
+    /* Similar appears to be true for GCN.  */
+    if (acc_get_device_type () == acc_device_radeon)
+      gangs = 3;
     int gangs_actual = gangs;
 #define WORKERS 3
     int workers_actual = WORKERS;
@@ -459,6 +488,13 @@ int main ()
 	  /* The GCC nvptx back end enforces vector_length (32).  */
 	  vectors_actual = 32;
 	}
+      else if (acc_on_device (acc_device_radeon))
+	{
+	  /* Temporary setting, until multiple workers are permitted.  */
+	  workers_actual = 1;
+	  /* See above comments about GCN vectors_actual.  */
+	  vectors_actual = 1;
+	}
       else
 	__builtin_abort ();
 #pragma acc loop gang reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c
index 2cb5b95..6570c64 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c
@@ -15,4 +15,22 @@ main (void)
   return 0;
 }
 
-/* { dg-final { scan-assembler-times "bar.sync" 0 } } */
+/* Todo: Boths bar.syncs can be removed.
+   Atm we generate this dead code inbetween forked and joining:
+
+                     mov.u32 %r28, %ntid.y;
+                     mov.u32 %r29, %tid.y;
+                     add.u32 %r30, %r29, %r29;
+                     setp.gt.s32     %r31, %r30, 19;
+             @%r31   bra     $L2;
+                     add.u32 %r25, %r28, %r28;
+                     mov.u32 %r24, %r30;
+     $L3:
+                     add.u32 %r24, %r24, %r25;
+                     setp.le.s32     %r33, %r24, 19;
+             @%r33   bra     $L3;
+     $L2:
+
+   so the loop is not recognized as empty loop (which we detect by seeing if
+   joining immediately follows forked).  */
+/* { dg-final { scan-assembler-times "bar.sync" 2 } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c
index e8a433f..d955d79 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c
@@ -21,4 +21,7 @@ main (void)
   return 0;
 }
 
-/* { dg-final { scan-assembler-times "bar.sync" 0 } } */
+/* Atm, %ntid.y is broadcast from one loop to the next, so there are 2 bar.syncs
+   for that (the other two are there for the same reason as in pr85381-2.c).
+   Todo: Recompute %ntid.y instead of broadcasting it. */
+/* { dg-final { scan-assembler-times "bar.sync" 4 } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c
index 6ba96b6..79cebf6 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c
@@ -1,31 +1,61 @@
-/* Verify that 'acc_unmap_data' unmaps even in presence of dynamic reference
-   counts.  */
+/* Verify that 'acc_unmap_data' unmaps even in presence of structured and
+   dynamic reference counts, but the device memory remains allocated.  */
 
 /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
 
 #include <assert.h>
 #include <stdlib.h>
+#include <string.h>
 #include <openacc.h>
 
 int
 main ()
 {
   const int N = 180;
-
-  char *h = (char *) malloc (N);
-  char *d = (char *) acc_malloc (N);
-  if (!d)
-    abort ();
-  acc_map_data (h, d, N);
-
-  char *d_ = (char *) acc_create (h + 3, N - 77);
-  assert (d_ == d + 3);
-
-  d_ = (char *) acc_create (h, N);
-  assert (d_ == d);
-
-  acc_unmap_data (h);
-  assert (!acc_is_present (h, N));
+  const int N_i = 537;
+  const int C = 37;
+
+  unsigned char *h = (unsigned char *) malloc (N);
+  assert (h);
+  unsigned char *d = (unsigned char *) acc_malloc (N);
+  assert (d);
+
+  for (int i = 0; i < N_i; ++i)
+    {
+      acc_map_data (h, d, N);
+      assert (acc_is_present (h, N));
+#pragma acc parallel present(h[0:N])
+      {
+	if (i == 0)
+	  memset (h, C, N);
+      }
+
+      unsigned char *d_ = (unsigned char *) acc_create (h + 3, N - 77);
+      assert (d_ == d + 3);
+
+#pragma acc data create(h[6:N - 44])
+      {
+	d_ = (unsigned char *) acc_create (h, N);
+	assert (d_ == d);
+
+#pragma acc enter data create(h[0:N])
+
+	assert (acc_is_present (h, N));
+	acc_unmap_data (h);
+	assert (!acc_is_present (h, N));
+      }
+
+      /* We can however still access the device memory.  */
+#pragma acc parallel loop deviceptr(d)
+      for (int j = 0; j < N; ++j)
+	d[j] += i * j;
+    }
+
+  acc_memcpy_from_device(h, d, N);
+  for (int j = 0; j < N; ++j)
+    assert (h[j] == ((C + N_i * (N_i - 1) / 2 * j) % 256));
+
+  acc_free (d);
 
   return 0;
 }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c
index a97e046..da13d84 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c
@@ -30,14 +30,18 @@ int main ()
   int ix;
   int exit = 0;
   int ondev = 0;
+  int gangsize, workersize, vectorsize;
 
   for (ix = 0; ix < N;ix++)
     ary[ix] = -1;
   
-#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) copy(ary) copy(ondev)
+#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) copy(ary) copy(ondev) copyout(gangsize, workersize, vectorsize)
   {
     ondev = acc_on_device (acc_device_not_host);
     gang (ary);
+    gangsize = __builtin_goacc_parlevel_size (GOMP_DIM_GANG);
+    workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER);
+    vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR);
   }
 
   for (ix = 0; ix < N; ix++)
@@ -45,11 +49,12 @@ int main ()
       int expected = ix;
       if(ondev)
 	{
-	  int chunk_size = (N + 32*32*32 - 1) / (32*32*32);
+	  int chunk_size = (N + gangsize * workersize * vectorsize - 1)
+			   / (gangsize * workersize * vectorsize);
 	  
-	  int g = ix / (chunk_size * 32 * 32);
-	  int w = ix / 32 % 32;
-	  int v = ix % 32;
+	  int g = ix / (chunk_size * vectorsize * workersize);
+	  int w = (ix / vectorsize) % workersize;
+	  int v = ix % vectorsize;
 
 	  expected = (g << 16) | (w << 8) | v;
 	}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c
index b1e3e3a..dd7bb6c 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c
@@ -30,14 +30,17 @@ int main ()
   int ix;
   int exit = 0;
   int ondev = 0;
+  int vectorsize;
 
   for (ix = 0; ix < N;ix++)
     ary[ix] = -1;
   
-#pragma acc parallel vector_length(32) copy(ary) copy(ondev)
+#pragma acc parallel vector_length(32) copy(ary) copy(ondev) \
+	    copyout(vectorsize)
   {
     ondev = acc_on_device (acc_device_not_host);
     vector (ary);
+    vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR);
   }
 
   for (ix = 0; ix < N; ix++)
@@ -47,7 +50,7 @@ int main ()
 	{
 	  int g = 0;
 	  int w = 0;
-	  int v = ix % 32;
+	  int v = ix % vectorsize;
 
 	  expected = (g << 16) | (w << 8) | v;
 	}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
index 81f1e03..acd9884 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
@@ -30,14 +30,17 @@ int main ()
   int ix;
   int exit = 0;
   int ondev = 0;
+  int workersize;
 
   for (ix = 0; ix < N;ix++)
     ary[ix] = -1;
   
-#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev)
+#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) \
+	    copyout(workersize)
   {
     ondev = acc_on_device (acc_device_not_host);
     worker (ary);
+    workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER);
   }
 
   for (ix = 0; ix < N; ix++)
@@ -46,7 +49,7 @@ int main ()
       if(ondev)
 	{
 	  int g = 0;
-	  int w = ix % 32;
+	  int w = ix % workersize;
 	  int v = 0;
 
 	  expected = (g << 16) | (w << 8) | v;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c
index 23dbc1a..73696e4 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c
@@ -30,14 +30,18 @@ int main ()
   int ix;
   int exit = 0;
   int ondev = 0;
+  int workersize, vectorsize;
 
   for (ix = 0; ix < N;ix++)
     ary[ix] = -1;
   
-#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev)
+#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) \
+	    copyout(workersize, vectorsize)
   {
     ondev = acc_on_device (acc_device_not_host);
     worker (ary);
+    workersize = __builtin_goacc_parlevel_size (GOMP_DIM_WORKER);
+    vectorsize = __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR);
   }
 
   for (ix = 0; ix < N; ix++)
@@ -46,8 +50,8 @@ int main ()
       if(ondev)
 	{
 	  int g = 0;
-	  int w = (ix / 32) % 32;
-	  int v = ix % 32;
+	  int w = (ix / vectorsize) % workersize;
+	  int v = ix % vectorsize;
 
 	  expected = (g << 16) | (w << 8) | v;
 	}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c
index 8862148..9769ee7 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c
@@ -2,8 +2,13 @@
 #include <openacc.h>
 #include <gomp-constants.h>
 
+#ifdef ACC_DEVICE_TYPE_radeon
+#define NUM_WORKERS 16
+#define NUM_VECTORS 1
+#else
 #define NUM_WORKERS 16
 #define NUM_VECTORS 32
+#endif
 #define WIDTH 64
 #define HEIGHT 32
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
new file mode 100644
index 0000000..543aaa15
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
@@ -0,0 +1,187 @@
+/* Test dynamic refcount and copy behavior of separate structure members.  */
+
+#include <assert.h>
+#include <stdbool.h>
+#include <openacc.h>
+
+struct s
+{
+  signed char a;
+  float b;
+};
+
+static void test(unsigned variant)
+{
+  struct s s = { .a = 73, .b = -22 };
+
+#pragma acc enter data copyin(s.a, s.b)
+  assert(acc_is_present(&s.a, sizeof s.a));
+  assert(acc_is_present(&s.b, sizeof s.b));
+
+  /* To verify that any following 'copyin' doesn't 'copyin' again.  */
+  s.a = -s.a;
+  s.b = -s.b;
+
+  if (variant & 4)
+    {
+      if (variant & 8)
+	{
+#pragma acc enter data copyin(s.b)
+	}
+      else
+	acc_copyin(&s.b, sizeof s.b);
+      assert(acc_is_present(&s.a, sizeof s.a));
+      assert(acc_is_present(&s.b, sizeof s.b));
+
+      if (variant & 16)
+	{
+#pragma acc enter data copyin(s.a)
+	}
+      else
+	acc_copyin(&s.a, sizeof s.a);
+      assert(acc_is_present(&s.a, sizeof s.a));
+      assert(acc_is_present(&s.b, sizeof s.b));
+
+      if (variant & 32)
+	{
+#pragma acc enter data copyin(s.a)
+	  acc_copyin(&s.b, sizeof s.b);
+#pragma acc enter data copyin(s.b)
+#pragma acc enter data copyin(s.b)
+	  acc_copyin(&s.a, sizeof s.a);
+	  acc_copyin(&s.a, sizeof s.a);
+	  acc_copyin(&s.a, sizeof s.a);
+	}
+      assert(acc_is_present(&s.a, sizeof s.a));
+      assert(acc_is_present(&s.b, sizeof s.b));
+    }
+
+#pragma acc parallel \
+  copy(s.a, s.b)
+  {
+#if ACC_MEM_SHARED
+    if (s.a++ != -73)
+      __builtin_abort();
+    if (s.b-- != 22)
+      __builtin_abort();
+#else
+    if (s.a++ != 73)
+      __builtin_abort();
+    if (s.b-- != -22)
+      __builtin_abort();
+#endif
+  }
+#if ACC_MEM_SHARED
+  assert(s.a == -72);
+  assert(s.b == 21);
+#else
+  assert(s.a == -73);
+  assert(s.b == 22);
+#endif
+
+  if (variant & 32)
+    {
+      if (variant & 1)
+	{
+#pragma acc exit data copyout(s.a) finalize
+	}
+      else
+	acc_copyout_finalize(&s.a, sizeof s.a);
+    }
+  else
+    {
+      if (variant & 1)
+	{
+#pragma acc exit data copyout(s.a)
+	}
+      else
+	acc_copyout(&s.a, sizeof s.a);
+      if (variant & 4)
+	{
+	  assert(acc_is_present(&s.a, sizeof s.a));
+	  assert(acc_is_present(&s.b, sizeof s.b));
+#if ACC_MEM_SHARED
+	  assert(s.a == -72);
+	  assert(s.b == 21);
+#else
+	  assert(s.a == -73);
+	  assert(s.b == 22);
+#endif
+	  if (variant & 1)
+	    {
+#pragma acc exit data copyout(s.a)
+	    }
+	  else
+	    acc_copyout(&s.a, sizeof s.a);
+	}
+    }
+#if ACC_MEM_SHARED
+  assert(acc_is_present(&s.a, sizeof s.a));
+  assert(acc_is_present(&s.b, sizeof s.b));
+  assert(s.a == -72);
+  assert(s.b == 21);
+#else
+  assert(!acc_is_present(&s.a, sizeof s.a));
+  assert(acc_is_present(&s.b, sizeof s.b));
+  assert(s.a == 74);
+  assert(s.b == 22);
+#endif
+
+  if (variant & 32)
+    {
+      if (variant & 2)
+	{
+#pragma acc exit data copyout(s.b) finalize
+	}
+      else
+	acc_copyout_finalize(&s.b, sizeof s.b);
+    }
+  else
+    {
+      if (variant & 2)
+	{
+#pragma acc exit data copyout(s.b)
+	}
+      else
+	acc_copyout(&s.b, sizeof s.b);
+      if (variant & 4)
+	{
+#if ACC_MEM_SHARED
+	  assert(acc_is_present(&s.a, sizeof s.a));
+	  assert(acc_is_present(&s.b, sizeof s.b));
+	  assert(s.a == -72);
+	  assert(s.b == 21);
+#else
+	  assert(!acc_is_present(&s.a, sizeof s.a));
+	  assert(acc_is_present(&s.b, sizeof s.b));
+	  assert(s.a == 74);
+	  assert(s.b == 22);
+#endif
+	  if (variant & 2)
+	    {
+#pragma acc exit data copyout(s.b)
+	    }
+	  else
+	    acc_copyout(&s.b, sizeof s.b);
+	}
+    }
+#if ACC_MEM_SHARED
+  assert(acc_is_present(&s.a, sizeof s.a));
+  assert(acc_is_present(&s.b, sizeof s.b));
+  assert(s.a == -72);
+  assert(s.b == 21);
+#else
+  assert(!acc_is_present(&s.a, sizeof s.a));
+  assert(!acc_is_present(&s.b, sizeof s.b));
+  assert(s.a == 74);
+  assert(s.b == -23);
+#endif
+}
+
+int main()
+{
+  for (unsigned variant = 0; variant < 64; ++variant)
+    test(variant);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
new file mode 100644
index 0000000..b86f1c9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
@@ -0,0 +1,38 @@
+#include <assert.h>
+
+struct str1 {
+  int a;
+  int b;
+};
+
+struct str2 {
+  int c;
+  int d;
+  struct str1 s;
+};
+
+int
+main (int argc, char *argv[])
+{
+  struct str2 t;
+
+  t.c = 1;
+  t.d = 2;
+  t.s.a = 3;
+  t.s.b = 4;
+
+  #pragma acc enter data copyin(t.s)
+
+  #pragma acc serial present(t.s) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
+  {
+    t.s.a = 5;
+    t.s.b = 6;
+  }
+
+  #pragma acc exit data copyout(t.s)
+
+  assert (t.s.a == 5);
+  assert (t.s.b == 6);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c
new file mode 100644
index 0000000..4dd8a3a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c
@@ -0,0 +1,44 @@
+#include <assert.h>
+#include <stdlib.h>
+
+struct str1 {
+  int a;
+  int b;
+  int *c;
+};
+
+#define N 1024
+
+int
+main (int argc, char *argv[])
+{
+  struct str1 s;
+
+  s.a = 1;
+  s.b = 2;
+  s.c = (int *) malloc (sizeof (int) * N);
+
+  for (int i = 0; i < N; i++)
+    s.c[i] = i + 10;
+
+  #pragma acc enter data copyin(s.a, s.b, s.c[0:N])
+
+  #pragma acc serial present(s.a, s.b, s.c[0:N]) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
+  {
+    s.a = 3;
+    s.b = 4;
+    for (int i = 0; i < N; i++)
+      s.c[i] = i + 20;
+  }
+
+  #pragma acc exit data copyout(s.a, s.b, s.c[0:N])
+
+  assert (s.a == 3);
+  assert (s.b == 4);
+  for (int i = 0; i < N; i++)
+    assert (s.c[i] == i + 20);
+
+  free (s.c);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c
new file mode 100644
index 0000000..8fa87777
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-1.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c
new file mode 100644
index 0000000..0d6b415
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c
@@ -0,0 +1,161 @@
+/* Test transitioning of data lifetimes between structured and dynamic.  */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define SIZE 1024
+
+void
+f1 (void)
+{
+  char *block1 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+  acc_copyin (block1, SIZE);
+  acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE])
+  {
+#ifdef OPENACC_API
+    acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+  }
+
+  assert (acc_is_present (block1, SIZE));
+
+#ifdef OPENACC_API
+  acc_copyout (block1, SIZE);
+  assert (acc_is_present (block1, SIZE));
+  acc_copyout (block1, SIZE);
+  assert (acc_is_present (block1, SIZE));
+  acc_copyout (block1, SIZE);
+  assert (!acc_is_present (block1, SIZE));
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+  assert (acc_is_present (block1, SIZE));
+#pragma acc exit data copyout(block1[0:SIZE])
+  assert (acc_is_present (block1, SIZE));
+#pragma acc exit data copyout(block1[0:SIZE])
+  assert (!acc_is_present (block1, SIZE));
+#endif
+
+  free (block1);
+}
+
+void
+f2 (void)
+{
+  char *block1 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+  acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE])
+  {
+#ifdef OPENACC_API
+    acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+    /* This should stay present until the end of the structured data
+       lifetime.  */
+    assert (acc_is_present (block1, SIZE));
+  }
+
+  assert (!acc_is_present (block1, SIZE));
+
+  free (block1);
+}
+
+void
+f3 (void)
+{
+  char *block1 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+  acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE])
+  {
+#ifdef OPENACC_API
+    acc_copyout (block1, SIZE);
+    acc_copyin (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+    assert (acc_is_present (block1, SIZE));
+  }
+
+  assert (acc_is_present (block1, SIZE));
+#ifdef OPENACC_API
+  acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+  assert (!acc_is_present (block1, SIZE));
+
+  free (block1);
+}
+
+void
+f4 (void)
+{
+  char *block1 = (char *) malloc (SIZE);
+  char *block2 = (char *) malloc (SIZE);
+  char *block3 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE], block3[0:SIZE])
+  {
+  /* The first copyin of block2 is the enclosing data region.  This
+     "enter data" should make it live beyond the end of this region.
+     This works, though the on-target copies of block1, block2 and block3
+     will stay allocated until block2 is unmapped because they are bound
+     together in a single target_mem_desc.  */
+#ifdef OPENACC_API
+    acc_copyin (block2, SIZE);
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+#endif
+  }
+
+  assert (!acc_is_present (block1, SIZE));
+  assert (acc_is_present (block2, SIZE));
+  assert (!acc_is_present (block3, SIZE));
+
+#ifdef OPENACC_API
+  acc_copyout (block2, SIZE);
+#else
+#pragma acc exit data copyout(block2[0:SIZE])
+#endif
+  assert (!acc_is_present (block2, SIZE));
+
+  free (block1);
+  free (block2);
+  free (block3);
+}
+
+int
+main (int argc, char *argv[])
+{
+  f1 ();
+  f2 ();
+  f3 ();
+  f4 ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c
new file mode 100644
index 0000000..365df8d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-2.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c
new file mode 100644
index 0000000..726942c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c
@@ -0,0 +1,166 @@
+/* Test nested dynamic/structured data mappings.  */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define SIZE 1024
+
+void
+f1 (void)
+{
+  char *block1 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE])
+  {
+#ifdef OPENACC_API
+    acc_copyin (block1, SIZE);
+    acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+  }
+
+  assert (!acc_is_present (block1, SIZE));
+
+  free (block1);
+}
+
+void
+f2 (void)
+{
+  char *block1 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+  acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE])
+  {
+  }
+
+#ifdef OPENACC_API
+  acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+
+  assert (!acc_is_present (block1, SIZE));
+
+  free (block1);
+}
+
+void
+f3 (void)
+{
+  char *block1 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE])
+  {
+#ifdef OPENACC_API
+    acc_copyin (block1, SIZE);
+    acc_copyin (block1, SIZE);
+    acc_copyout (block1, SIZE);
+    acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+  }
+
+  assert (!acc_is_present (block1, SIZE));
+
+  free (block1);
+}
+
+void
+f4 (void)
+{
+  char *block1 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE])
+  {
+#ifdef OPENACC_API
+    acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE])
+    {
+#ifdef OPENACC_API
+      acc_copyin (block1, SIZE);
+      acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+    }
+
+#ifdef OPENACC_API
+  acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+  }
+
+  assert (!acc_is_present (block1, SIZE));
+
+  free (block1);
+}
+
+void
+f5 (void)
+{
+  char *block1 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+  acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE])
+  {
+#ifdef OPENACC_API
+    acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+#pragma acc data copy(block1[0:SIZE])
+    {
+    }
+#ifdef OPENACC_API
+    acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+  }
+#ifdef OPENACC_API
+  acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+
+  assert (!acc_is_present (block1, SIZE));
+
+  free (block1);
+}
+
+int
+main (int argc, char *argv[])
+{
+  f1 ();
+  f2 ();
+  f3 ();
+  f4 ();
+  f5 ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c
new file mode 100644
index 0000000..469b35b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-3.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c
new file mode 100644
index 0000000..c13f3c5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c
@@ -0,0 +1,183 @@
+/* Test nested dynamic/structured data mappings (multiple blocks on data
+   regions).  */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define SIZE 1024
+
+void
+f1 (void)
+{
+  char *block1 = (char *) malloc (SIZE);
+  char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+  {
+#ifdef OPENACC_API
+    acc_copyin (block1, SIZE);
+    acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+  }
+
+  assert (!acc_is_present (block1, SIZE));
+  assert (!acc_is_present (block2, SIZE));
+
+  free (block1);
+  free (block2);
+}
+
+void
+f2 (void)
+{
+  char *block1 = (char *) malloc (SIZE);
+  char *block2 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+  acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+  {
+  }
+
+#ifdef OPENACC_API
+  acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+
+  assert (!acc_is_present (block1, SIZE));
+  assert (!acc_is_present (block2, SIZE));
+
+  free (block1);
+  free (block2);
+}
+
+void
+f3 (void)
+{
+  char *block1 = (char *) malloc (SIZE);
+  char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+  {
+#ifdef OPENACC_API
+    acc_copyin (block1, SIZE);
+    acc_copyin (block2, SIZE);
+    acc_copyout (block2, SIZE);
+    acc_copyout (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc enter data copyin(block2[0:SIZE])
+#pragma acc exit data copyout(block2[0:SIZE])
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+  }
+
+  assert (!acc_is_present (block1, SIZE));
+  assert (!acc_is_present (block2, SIZE));
+
+  free (block1);
+  free (block2);
+}
+
+void
+f4 (void)
+{
+  char *block1 = (char *) malloc (SIZE);
+  char *block2 = (char *) malloc (SIZE);
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+  {
+#ifdef OPENACC_API
+    acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+    {
+#ifdef OPENACC_API
+      acc_copyin (block2, SIZE);
+      acc_copyout (block2, SIZE);
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+#pragma acc exit data copyout(block2[0:SIZE])
+#endif
+    }
+#ifdef OPENACC_API
+  acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+  }
+
+  assert (!acc_is_present (block1, SIZE));
+  assert (!acc_is_present (block2, SIZE));
+
+  free (block1);
+  free (block2);
+}
+
+void
+f5 (void)
+{
+  char *block1 = (char *) malloc (SIZE);
+  char *block2 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+  acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+  {
+#ifdef OPENACC_API
+    acc_copyin (block2, SIZE);
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+    {
+    }
+#ifdef OPENACC_API
+    acc_copyout (block2, SIZE);
+#else
+#pragma acc exit data copyout(block2[0:SIZE])
+#endif
+  }
+
+#ifdef OPENACC_API
+  acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+
+  assert (!acc_is_present (block1, SIZE));
+  assert (!acc_is_present (block2, SIZE));
+
+  free (block1);
+  free (block2);
+}
+
+int
+main (int argc, char *argv[])
+{
+  f1 ();
+  f2 ();
+  f3 ();
+  f4 ();
+  f5 ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c
new file mode 100644
index 0000000..8e88b97
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-4.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c
new file mode 100644
index 0000000..e9a6510
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c
@@ -0,0 +1,64 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define SIZE 1024
+
+int
+main (int argc, char *argv[])
+{
+  char *block1 = (char *) malloc (SIZE);
+  char *block2 = (char *) malloc (SIZE);
+  char *block3 = (char *) malloc (SIZE);
+
+  /* Doing this twice ensures that we have a non-zero virtual refcount.  Make
+     sure that works too.  */
+#ifdef OPENACC_API
+  acc_copyin (block1, SIZE);
+  acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE], block3[0:SIZE])
+  {
+  /* The first copyin of block2 is the enclosing data region.  This
+     "enter data" should make it live beyond the end of this region.  */
+#ifdef OPENACC_API
+    acc_copyin (block2, SIZE);
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+#endif
+  }
+
+  assert (acc_is_present (block1, SIZE));
+  assert (acc_is_present (block2, SIZE));
+  assert (!acc_is_present (block3, SIZE));
+
+#ifdef OPENACC_API
+  acc_copyout (block1, SIZE);
+  assert (acc_is_present (block1, SIZE));
+  acc_copyout (block1, SIZE);
+  assert (!acc_is_present (block1, SIZE));
+
+  acc_copyout (block2, SIZE);
+  assert (!acc_is_present (block2, SIZE));
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+  assert (acc_is_present (block1, SIZE));
+#pragma acc exit data copyout(block1[0:SIZE])
+  assert (!acc_is_present (block1, SIZE));
+
+#pragma acc exit data copyout(block2[0:SIZE])
+  assert (!acc_is_present (block2, SIZE));
+#endif
+
+  free (block1);
+  free (block2);
+  free (block3);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c
new file mode 100644
index 0000000..59ef562
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-5.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c
new file mode 100644
index 0000000..9807076
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c
@@ -0,0 +1,56 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define SIZE 1024
+
+int
+main (int argc, char *argv[])
+{
+  char *block1 = (char *) malloc (SIZE);
+  char *block2 = (char *) malloc (SIZE);
+  char *block3 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+  acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE], block3[0:SIZE])
+  {
+  /* The first copyin of block2 is the enclosing data region.  This
+     "enter data" should make it live beyond the end of this region.  */
+#ifdef OPENACC_API
+    acc_copyin (block2, SIZE);
+#else
+#pragma acc enter data copyin(block2[0:SIZE])
+#endif
+  }
+
+  assert (acc_is_present (block1, SIZE));
+  assert (acc_is_present (block2, SIZE));
+  assert (!acc_is_present (block3, SIZE));
+
+#ifdef OPENACC_API
+  acc_copyout (block1, SIZE);
+  assert (!acc_is_present (block1, SIZE));
+
+  acc_copyout (block2, SIZE);
+  assert (!acc_is_present (block2, SIZE));
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+  assert (!acc_is_present (block1, SIZE));
+
+#pragma acc exit data copyout(block2[0:SIZE])
+  assert (!acc_is_present (block2, SIZE));
+#endif
+
+  free (block1);
+  free (block2);
+  free (block3);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c
new file mode 100644
index 0000000..0401f73
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-6.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c
new file mode 100644
index 0000000..9250b4a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c
@@ -0,0 +1,43 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define SIZE 1024
+
+int
+main (int argc, char *argv[])
+{
+  char *block1 = (char *) malloc (SIZE);
+  char *block2 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+  acc_copyin (block1, SIZE);
+  acc_copyin (block2, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE], block2[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+  {
+#ifdef OPENACC_API
+    acc_copyout (block1, SIZE);
+    acc_copyout (block2, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE], block2[0:SIZE])
+#endif
+    /* These should stay present until the end of the structured data
+       lifetime.  */
+    assert (acc_is_present (block1, SIZE));
+    assert (acc_is_present (block2, SIZE));
+  }
+
+  assert (!acc_is_present (block1, SIZE));
+  assert (!acc_is_present (block2, SIZE));
+
+  free (block1);
+  free (block2);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c
new file mode 100644
index 0000000..07caefb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-7.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c
new file mode 100644
index 0000000..52e8d4c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c
@@ -0,0 +1,44 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define SIZE 1024
+
+int
+main (int argc, char *argv[])
+{
+  char *block1 = (char *) malloc (SIZE);
+  char *block2 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+  acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+  {
+/* We can't attach the dynamic data mapping's (block1) target_mem_desc to the
+   enclosing structured data region here, because that region maps block2
+   also.  */
+#ifdef OPENACC_API
+    acc_copyout (block1, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#endif
+    /* These should stay present until the end of the structured data
+       lifetime.  */
+    assert (acc_is_present (block1, SIZE));
+    assert (acc_is_present (block2, SIZE));
+  }
+
+  assert (!acc_is_present (block1, SIZE));
+  assert (!acc_is_present (block2, SIZE));
+
+  free (block1);
+  free (block2);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c
new file mode 100644
index 0000000..1c2479a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c
@@ -0,0 +1,3 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+/* { dg-additional-options "-DOPENACC_API" } */
+#include "structured-dynamic-lifetimes-8.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c
new file mode 100644
index 0000000..919ee02
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c
@@ -0,0 +1,47 @@
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <openacc.h>
+#include <assert.h>
+#include <stdlib.h>
+
+#define SIZE 1024
+
+int
+main (int argc, char *argv[])
+{
+  char *block1 = (char *) malloc (SIZE);
+  char *block2 = (char *) malloc (SIZE);
+
+#ifdef OPENACC_API
+  acc_copyin (block1, SIZE);
+#else
+#pragma acc enter data copyin(block1[0:SIZE])
+#endif
+
+#pragma acc data copy(block1[0:SIZE], block2[0:SIZE])
+  {
+#ifdef OPENACC_API
+    acc_copyout (block1, SIZE);
+    acc_copyin (block2, SIZE);
+#else
+#pragma acc exit data copyout(block1[0:SIZE])
+#pragma acc enter data copyin(block2[0:SIZE])
+#endif
+    assert (acc_is_present (block1, SIZE));
+    assert (acc_is_present (block2, SIZE));
+  }
+
+  assert (!acc_is_present (block1, SIZE));
+  assert (acc_is_present (block2, SIZE));
+#ifdef OPENACC_API
+  acc_copyout (block2, SIZE);
+#else
+#pragma acc exit data copyout(block2[0:SIZE])
+#endif
+  assert (!acc_is_present (block2, SIZE));
+
+  free (block1);
+  free (block2);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c
index c019fe5..57579171 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c
@@ -1,5 +1,5 @@
 /* AMD GCN does not use 32-lane vectors, so the expected use counts mismatch.
-   { dg-skip-if "unsuitable dimensions" { openacc_amdgcn_accel_selected } { "*" } { "" } } */
+   { dg-skip-if "unsuitable dimensions" { openacc_radeon_accel_selected } { "*" } { "" } } */
 
 /* { dg-additional-options "-fopenacc-dim=32" } */
 
diff --git a/libgomp/testsuite/libgomp.oacc-c/c.exp b/libgomp/testsuite/libgomp.oacc-c/c.exp
index 7f13242..48cbc98 100644
--- a/libgomp/testsuite/libgomp.oacc-c/c.exp
+++ b/libgomp/testsuite/libgomp.oacc-c/c.exp
@@ -51,15 +51,6 @@ foreach offload_target [concat [split $offload_targets ","] "disable"] {
 	    unsupported "$subdir $offload_target offloading"
 	    continue
 	}
-	gcn {
-	    if { ![check_effective_target_openacc_amdgcn_accel_present] } {
-		# Don't bother; execution testing is going to FAIL.
-		untested "$subdir $offload_target offloading: supported, but hardware not accessible"
-		continue
-	    }
-
-	    set acc_mem_shared 0
-	}
 	host {
 	    set acc_mem_shared 1
 	}
@@ -78,6 +69,15 @@ foreach offload_target [concat [split $offload_targets ","] "disable"] {
 
 	    set acc_mem_shared 0
 	}
+	radeon {
+	    if { ![check_effective_target_openacc_radeon_accel_present] } {
+		# Don't bother; execution testing is going to FAIL.
+		untested "$subdir $offload_target offloading: supported, but hardware not accessible"
+		continue
+	    }
+
+	    set acc_mem_shared 0
+	}
 	default {
 	    error "Unknown OpenACC device type: $openacc_device_type (offload target: $offload_target)"
 	}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/acc_get_property.f90 b/libgomp/testsuite/libgomp.oacc-fortran/acc_get_property.f90
index ce69547..1af7cc3 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/acc_get_property.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/acc_get_property.f90
@@ -3,8 +3,6 @@
 ! of all device types mentioned in the OpenACC standard.
 !
 ! See also acc_get_property.c
-! { dg-do run { target { { ! { openacc_host_selected } } && { ! { openacc_amdgcn_accel_selected } } } } }
-! FIXME: This test does not work with the GCN implementation stub yet.
 
 program test
   use openacc
@@ -28,13 +26,14 @@ end program test
 ! and do basic device independent validation.
 subroutine print_device_properties (device_type)
   use openacc
+  use iso_c_binding, only: c_size_t
   implicit none
 
   integer, intent(in) :: device_type
 
   integer :: device_count
   integer :: device
-  integer(acc_device_property) :: v
+  integer(c_size_t) :: v
   character*256 :: s
 
   device_count = acc_get_num_devices(device_type)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/atomic_capture-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/atomic_capture-1.f90
index 5a4a1e0..536b3f0 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/atomic_capture-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/atomic_capture-1.f90
@@ -275,8 +275,9 @@ program main
   if (ltmp .neqv. .not. lexp) STOP 33
   if (lgot .neqv. lexp) STOP 34
 
-  igot = 1
+  igot = 0
   iexp = N
+  iarr = -42
 
   !$acc parallel loop copy (igot, itmp)
     do i = 1, N
@@ -287,13 +288,24 @@ program main
     end do
   !$acc end parallel loop
 
+  if (igot /= N) stop 107
+  itmp = 0
+  do i = 1, N
+     if (iarr(i) == 0) then
+       itmp = i
+       exit
+     end if
+  end do
+  ! At most one iarr element can be 0.
   do i = 1, N
-     if (.not. (1 <= iarr(i) .and. iarr(i) < iexp)) STOP 35
+     if ((iarr(i) == 0 .and. i /= itmp) &
+         .or. iarr(i) < 0 .or. iarr(i) >= N) STOP 35
   end do
   if (igot /= iexp) STOP 36
 
-  igot = N
+  igot = N + 1
   iexp = 1
+  iarr = -42
 
   !$acc parallel loop copy (igot, itmp)
     do i = 1, N
@@ -304,8 +316,18 @@ program main
     end do
   !$acc end parallel loop
 
+  if (igot /= 1) stop 108
+  itmp = N + 1
+  ! At most one iarr element can be N+1.
+  do i = 1, N
+     if (iarr(i) == N + 1) then
+       itmp = i
+       exit
+     end if
+  end do
   do i = 1, N
-     if (.not. (iarr(i) == 1 .or. iarr(i) == N)) STOP 37
+     if ((iarr(i) == N + 1 .and. i /= itmp) &
+         .or. iarr(i) <= 0 .or. iarr(i) > N + 1) STOP 37
   end do
   if (igot /= iexp) STOP 38
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/classtypes-1.f95 b/libgomp/testsuite/libgomp.oacc-fortran/classtypes-1.f95
index f16f42f..c5f0fff 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/classtypes-1.f95
+++ b/libgomp/testsuite/libgomp.oacc-fortran/classtypes-1.f95
@@ -31,7 +31,8 @@ program main
     myvar%p%p(i) = -1.0
   end do
 
-!$acc enter data copyin(myvar, myvar%p) create(myvar%p%p)
+!$acc enter data copyin(myvar)
+!$acc enter data copyin(myvar%p) create(myvar%p%p)
 
 !$acc parallel loop present(myvar%p%p)
   do i=1,100
@@ -39,7 +40,8 @@ program main
   end do
 !$acc end parallel loop
 
-!$acc exit data copyout(myvar%p%p) delete(myvar, myvar%p)
+!$acc exit data copyout(myvar%p%p) delete(myvar%p)
+!$acc exit data delete(myvar)
 
   do i=1,100
     if (myvar%p%p(i) .ne. i * 2) stop 1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90
deleted file mode 100644
index 3593661..0000000
--- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90
+++ /dev/null
@@ -1,33 +0,0 @@
-! { dg-do run }
-
-! Test of attach/detach with "acc data", two clauses at once.
-
-program dtype
-  implicit none
-  integer, parameter :: n = 512
-  type mytype
-    integer, allocatable :: a(:)
-  end type mytype
-  integer i
-
-  type(mytype) :: var
-
-  allocate(var%a(1:n))
-
-!$acc data copy(var) copy(var%a)
-
-!$acc parallel loop
-  do i = 1,n
-    var%a(i) = i
-  end do
-!$acc end parallel loop
-
-!$acc end data
-
-  do i = 1,n
-    if (i .ne. var%a(i)) stop 1
-  end do
-
-  deallocate(var%a)
-
-end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90
index 667d944..edb6b8d 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90
@@ -16,12 +16,14 @@ program dtype
   allocate(var%a(1:n))
   allocate(var%b(1:n))
 
-!$acc parallel loop copy(var) copy(var%a(1:n)) copy(var%b(1:n))
+!$acc data copy(var)
+!$acc parallel loop copy(var%a(1:n)) copy(var%b(1:n))
   do i = 1,n
     var%a(i) = i
     var%b(i) = i
   end do
 !$acc end parallel loop
+!$acc end data
 
   do i = 1,n
     if (i .ne. var%a(i)) stop 1
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90
new file mode 100644
index 0000000..ed4f10e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90
@@ -0,0 +1,8 @@
+! { dg-do run }
+
+/* Nullify the 'finalize' clause, which disturbs reference counting.  */
+#define finalize
+#include "deep-copy-6.f90"
+
+! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
+! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
index 12910d0..5837a40 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
@@ -3,6 +3,7 @@
 ! Test of attachment counters and finalize.
 
 program dtype
+  use openacc
   implicit none
   integer, parameter :: n = 512
   type mytype
@@ -36,7 +37,23 @@ program dtype
   end do
 !$acc end parallel loop
 
+  if (.not. acc_is_present(var%a(5:n - 5))) stop 11
+  if (.not. acc_is_present(var%b(5:n - 5))) stop 12
+  if (.not. acc_is_present(var)) stop 13
+  print *, "CheCKpOInT1"
+  ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
 !$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5)) finalize
+  !TODO     goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
+  !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+  !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
+  !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
+  print *, "CheCKpOInT2"
+  ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
+  if (acc_get_device_type() .ne. acc_device_host) then
+     if (acc_is_present(var%a(5:n - 5))) stop 21
+     if (acc_is_present(var%b(5:n - 5))) stop 22
+  end if
+  if (.not. acc_is_present(var)) stop 23
 
 !$acc end data
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f
index e7358f4..de72774 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-1.f
@@ -3,6 +3,10 @@
       PROGRAM MAIN
       IMPLICIT NONE
 
+! Initialize before the checkpoint, in case this produces any output.
+!$ACC PARALLEL
+!$ACC END PARALLEL
+
       PRINT *, "CheCKpOInT"
 !$ACC PARALLEL
       ERROR STOP
@@ -17,7 +21,7 @@
 ! In gfortran's main program, libfortran's set_options is called - which sets
 ! compiler_options.backtrace = 1 by default.  For an offload libgfortran, this
 ! is never called and, hence, "Error termination." is never printed.  Thus:
-! { dg-output "Error termination.*" { target { ! { openacc_nvidia_accel_selected || openacc_amdgcn_accel_selected } } } }
+! { dg-output "Error termination.*" { target { ! { openacc_nvidia_accel_selected || openacc_radeon_accel_selected } } } }
 !
 ! PR85463:
 ! { dg-output "libgomp: cuStreamSynchronize error.*" { target openacc_nvidia_accel_selected } }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f
index fca1d96..475c9cb 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-2.f
@@ -3,6 +3,10 @@
       PROGRAM MAIN
       IMPLICIT NONE
 
+! Initialize before the checkpoint, in case this produces any output.
+!$ACC PARALLEL
+!$ACC END PARALLEL
+
       PRINT *, "CheCKpOInT"
 !$ACC PARALLEL
       ERROR STOP 35
@@ -17,7 +21,7 @@
 ! In gfortran's main program, libfortran's set_options is called - which sets
 ! compiler_options.backtrace = 1 by default.  For an offload libgfortran, this
 ! is never called and, hence, "Error termination." is never printed.  Thus:
-! { dg-output "Error termination.*" { target { ! { openacc_nvidia_accel_selected || openacc_amdgcn_accel_selected } } } }
+! { dg-output "Error termination.*" { target { ! { openacc_nvidia_accel_selected || openacc_radeon_accel_selected } } } }
 !
 ! PR85463:
 ! { dg-output "libgomp: cuStreamSynchronize error.*" { target openacc_nvidia_accel_selected } }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f
index 2ae0b0d..ab63444 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/error_stop-3.f
@@ -3,6 +3,10 @@
       PROGRAM MAIN
       IMPLICIT NONE
 
+! Initialize before the checkpoint, in case this produces any output.
+!$ACC PARALLEL
+!$ACC END PARALLEL
+
       PRINT *, "CheCKpOInT"
 !$ACC PARALLEL
       ERROR STOP "SiGN"
@@ -17,7 +21,7 @@
 ! In gfortran's main program, libfortran's set_options is called - which sets
 ! compiler_options.backtrace = 1 by default.  For an offload libgfortran, this
 ! is never called and, hence, "Error termination." is never printed.  Thus:
-! { dg-output "Error termination.*" { target { ! { openacc_nvidia_accel_selected || openacc_amdgcn_accel_selected } } } }
+! { dg-output "Error termination.*" { target { ! { openacc_nvidia_accel_selected || openacc_radeon_accel_selected } } } }
 !
 ! PR85463:
 ! { dg-output "libgomp: cuStreamSynchronize error.*" { target openacc_nvidia_accel_selected } }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
index 60f0889..d607903 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
+++ b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
@@ -82,8 +82,11 @@ if { $lang_test_file_found } {
 		unsupported "$subdir $offload_target offloading"
 		continue
 	    }
-	    gcn {
-		if { ![check_effective_target_openacc_amdgcn_accel_present] } {
+	    host {
+		set acc_mem_shared 1
+	    }
+	    nvidia {
+		if { ![check_effective_target_openacc_nvidia_accel_present] } {
 		    # Don't bother; execution testing is going to FAIL.
 		    untested "$subdir $offload_target offloading: supported, but hardware not accessible"
 		    continue
@@ -91,11 +94,8 @@ if { $lang_test_file_found } {
 
 		set acc_mem_shared 0
 	    }
-	    host {
-		set acc_mem_shared 1
-	    }
-	    nvidia {
-		if { ![check_effective_target_openacc_nvidia_accel_present] } {
+	    radeon {
+		if { ![check_effective_target_openacc_radeon_accel_present] } {
 		    # Don't bother; execution testing is going to FAIL.
 		    untested "$subdir $offload_target offloading: supported, but hardware not accessible"
 		    continue
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90 b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
new file mode 100644
index 0000000..483ac3f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
@@ -0,0 +1,92 @@
+! { dg-do run }
+!
+! Test if, if_present clauses on host_data construct.
+!
+! Fortran variant of 'libgomp.oacc-c-c++-common/host_data-7.c'.
+!
+program main
+  use iso_c_binding
+  implicit none
+  real, target :: var, arr(100)
+  integer(c_intptr_t) :: host_p, host_parr
+  host_p = transfer(c_loc(var), host_p)
+  host_parr = transfer(c_loc(arr), host_parr)
+  call foo (var, arr, host_p, host_parr, .false.)
+  call foo (var, arr, host_p, host_parr, .true.)
+
+contains
+
+subroutine foo (p2, parr, host_p, host_parr, cond)
+  use openacc
+  implicit none
+  real, target, intent(in) :: parr(:), p2
+  integer(c_intptr_t), value, intent(in) :: host_p, host_parr
+  logical, value, intent(in) :: cond
+  real, pointer :: p
+  p => p2
+
+  if (host_p /= transfer(c_loc(p), host_p)) stop 1
+  if (host_parr /= transfer(c_loc(parr), host_parr)) stop 2
+#if !ACC_MEM_SHARED
+  if (acc_is_present(p, c_sizeof(p))) stop 3
+  if (acc_is_present(parr, 1)) stop 4
+#endif
+  
+  !$acc data copyin(host_p, host_parr)
+#if !ACC_MEM_SHARED
+    if (acc_is_present(p, c_sizeof(p))) stop 5
+    if (acc_is_present(parr, 1)) stop 6
+#endif
+    !$acc host_data use_device(p, parr) if_present
+      ! not mapped yet, so it will be equal to the host pointer.
+      if (transfer(c_loc(p), host_p) /= host_p) stop 7
+      if (transfer(c_loc(parr), host_parr) /= host_parr) stop 8
+    !$acc end host_data
+#if !ACC_MEM_SHARED
+    if (acc_is_present(p, c_sizeof(p))) stop 9
+    if (acc_is_present(parr, 1)) stop 10
+#endif
+
+    !$acc data copy(p, parr)
+      if (.not. acc_is_present(p, c_sizeof(p))) stop 11
+      if (.not. acc_is_present(parr, 1)) stop 12
+      ! Not inside a host_data construct, so still the host pointer.
+      if (transfer(c_loc(p), host_p) /= host_p) stop 13
+      if (transfer(c_loc(parr), host_parr) /= host_parr) stop 14
+      
+      !$acc host_data use_device(p, parr)
+#if ACC_MEM_SHARED
+        if (transfer(c_loc(p), host_p) /= host_p) stop 15
+        if (transfer(c_loc(parr), host_parr) /= host_parr) stop 16
+#else
+        ! The device address is different from host address.
+        if (transfer(c_loc(p), host_p) == host_p) stop 17
+        if (transfer(c_loc(parr), host_parr) == host_parr) stop 18
+#endif
+      !$acc end host_data
+
+      !$acc host_data use_device(p, parr) if_present
+#if ACC_MEM_SHARED
+        if (transfer(c_loc(p), host_p) /= host_p) stop 19
+        if (transfer(c_loc(parr), host_parr) /= host_parr) stop 20
+#else
+        ! is present now, so this is the same as above.
+        if (transfer(c_loc(p), host_p) == host_p) stop 21
+        if (transfer(c_loc(parr), host_parr) == host_parr) stop 22
+#endif
+      !$acc end host_data
+
+      !$acc host_data use_device(p, parr) if(cond)
+#if ACC_MEM_SHARED
+        if (transfer(c_loc(p), host_p) /= host_p) stop 23
+        if (transfer(c_loc(parr), host_parr) /= host_parr) stop 24
+#else
+        ! is the device pointer iff cond is true.
+        if ((transfer(c_loc(p), host_p) /= host_p) .neqv. cond) stop 25
+        if ((transfer(c_loc(parr), host_parr) /= host_parr) .neqv. cond) stop 26
+#endif
+      !$acc end host_data
+    !$acc end data
+  !$acc end data
+end subroutine foo
+end
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90
new file mode 100644
index 0000000..445cbab
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90
@@ -0,0 +1,42 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+! Adapted from 'libgomp.oacc-fortran/deep-copy-6.f90'.
+
+program main
+  use openacc
+  implicit none
+  integer, parameter :: n = 512
+  type mytype
+    integer, allocatable :: a(:)
+  end type mytype
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+
+  !$acc data create(var)
+
+  !$acc enter data create(var%a)
+
+  if (.not. acc_is_present(var%a)) stop 1
+  if (.not. acc_is_present(var)) stop 2
+
+  print *, "CheCKpOInT1"
+  ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
+  !$acc exit data delete(var%a) finalize
+  !TODO     goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
+  !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+  !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
+  !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
+  print *, "CheCKpOInT2"
+  ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
+  if (acc_is_present(var%a)) stop 3
+  if (.not. acc_is_present(var)) stop 4
+
+  !$acc end data
+  if (acc_is_present(var%a)) stop 5
+  if (acc_is_present(var)) stop 6
+
+  deallocate(var%a)
+
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90
new file mode 100644
index 0000000..7b206ac
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90
@@ -0,0 +1,9 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+/* Nullify the 'finalize' clause, which disturbs reference counting.  */
+#define finalize
+#include "mdc-refcount-1-1-1.f90"
+
+! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
+! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90
new file mode 100644
index 0000000..8554534
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90
@@ -0,0 +1,44 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+! Adapted from 'libgomp.oacc-fortran/mdc-refcount-1-1-1.f90'.
+
+program main
+  use openacc
+  implicit none
+  integer, parameter :: n = 512
+  type mytype
+    integer, allocatable :: a(:)
+  end type mytype
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+
+  !$acc data create(var)
+
+  call acc_create(var%a)
+  ! After mapping via runtime API call, separately trigger attach action; see <https://github.com/OpenACC/openacc-spec/issues/301>.
+  !$acc enter data attach(var%a)
+
+  if (.not. acc_is_present(var%a)) stop 1
+  if (.not. acc_is_present(var)) stop 2
+
+  print *, "CheCKpOInT1"
+  ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
+  !$acc exit data delete(var%a) finalize
+  !TODO     goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
+  !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+  !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
+  !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
+  print *, "CheCKpOInT2"
+  ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
+  if (acc_is_present(var%a)) stop 3
+  if (.not. acc_is_present(var)) stop 4
+
+  !$acc end data
+  if (acc_is_present(var%a)) stop 5
+  if (acc_is_present(var)) stop 6
+
+  deallocate(var%a)
+
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90
new file mode 100644
index 0000000..8e696cc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90
@@ -0,0 +1,44 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+! Copy of 'libgomp.oacc-fortran/mdc-refcount-1-2-1.f90', without 'finalize' clause.
+
+program main
+  use openacc
+  implicit none
+  integer, parameter :: n = 512
+  type mytype
+    integer, allocatable :: a(:)
+  end type mytype
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+
+  !$acc data create(var)
+
+  call acc_create(var%a)
+  ! After mapping via runtime API call, separately trigger attach action; see <https://github.com/OpenACC/openacc-spec/issues/301>.
+  !$acc enter data attach(var%a)
+
+  if (.not. acc_is_present(var%a)) stop 1
+  if (.not. acc_is_present(var)) stop 2
+
+  print *, "CheCKpOInT1"
+  ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
+  !$acc exit data delete(var%a)
+  !TODO     goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
+  !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+  !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
+  !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
+  print *, "CheCKpOInT2"
+  ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
+  if (acc_is_present(var%a)) stop 3
+  if (.not. acc_is_present(var)) stop 4
+
+  !$acc end data
+  if (acc_is_present(var%a)) stop 5
+  if (acc_is_present(var)) stop 6
+
+  deallocate(var%a)
+
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90
new file mode 100644
index 0000000..070a6f8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90
@@ -0,0 +1,45 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+! Adapted from 'libgomp.oacc-fortran/mdc-refcount-1-2-1.f90'.
+
+program main
+  use openacc
+  implicit none
+  integer, parameter :: n = 512
+  type mytype
+    integer, allocatable :: a(:)
+  end type mytype
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+
+  !$acc data create(var)
+
+  call acc_create(var%a)
+  ! After mapping via runtime API call, separately trigger attach action; see <https://github.com/OpenACC/openacc-spec/issues/301>.
+  !$acc enter data attach(var%a)
+
+  if (.not. acc_is_present(var%a)) stop 1
+  if (.not. acc_is_present(var)) stop 2
+
+  !$acc exit data detach(var%a)
+  print *, "CheCKpOInT1"
+  ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
+  !$acc exit data delete(var%a) finalize
+  !TODO     goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
+  !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+  !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
+  !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
+  print *, "CheCKpOInT2"
+  ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
+  if (acc_is_present(var%a)) stop 3
+  if (.not. acc_is_present(var)) stop 4
+
+  !$acc end data
+  if (acc_is_present(var%a)) stop 5
+  if (acc_is_present(var)) stop 6
+
+  deallocate(var%a)
+
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90
new file mode 100644
index 0000000..3c4bbda
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90
@@ -0,0 +1,44 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+! Copy of 'libgomp.oacc-fortran/mdc-refcount-1-3-1.f90', without 'finalize' clause.
+
+program main
+  use openacc
+  implicit none
+  integer, parameter :: n = 512
+  type mytype
+    integer, allocatable :: a(:)
+  end type mytype
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+
+  !$acc data create(var)
+
+  call acc_create(var%a)
+  ! After mapping via runtime API call, separately trigger attach action; see <https://github.com/OpenACC/openacc-spec/issues/301>.
+  !$acc enter data attach(var%a)
+
+  if (.not. acc_is_present(var%a)) stop 1
+  if (.not. acc_is_present(var)) stop 2
+
+  !$acc exit data detach(var%a)
+  print *, "CheCKpOInT1"
+  ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
+  !$acc exit data delete(var%a)
+  !TODO { dg-output "(\n|\r\n|\r)libgomp: attach count underflow(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+  !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
+  !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
+  print *, "CheCKpOInT2"
+  ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
+  if (acc_is_present(var%a)) stop 3
+  if (.not. acc_is_present(var)) stop 4
+
+  !$acc end data
+  if (acc_is_present(var%a)) stop 5
+  if (acc_is_present(var)) stop 6
+
+  deallocate(var%a)
+
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90
new file mode 100644
index 0000000..b22e411
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90
@@ -0,0 +1,45 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+! Adapted from 'libgomp.oacc-fortran/mdc-refcount-1-3-1.f90'.
+
+program main
+  use openacc
+  implicit none
+  integer, parameter :: n = 512
+  type mytype
+    integer, allocatable :: a(:)
+  end type mytype
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+
+  !$acc data create(var)
+
+  call acc_create(var%a)
+  ! After mapping via runtime API call, separately trigger attach action; see <https://github.com/OpenACC/openacc-spec/issues/301>.
+  !$acc enter data attach(var%a)
+
+  if (.not. acc_is_present(var%a)) stop 1
+  if (.not. acc_is_present(var)) stop 2
+
+  print *, "CheCKpOInT1"
+  ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
+  !$acc exit data detach(var%a) finalize
+  !TODO     goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
+  !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+  !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
+  !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
+  print *, "CheCKpOInT2"
+  ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
+  !$acc exit data delete(var%a)
+  if (acc_is_present(var%a)) stop 3
+  if (.not. acc_is_present(var)) stop 4
+
+  !$acc end data
+  if (acc_is_present(var%a)) stop 5
+  if (acc_is_present(var)) stop 6
+
+  deallocate(var%a)
+
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90
new file mode 100644
index 0000000..476cd5c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90
@@ -0,0 +1,44 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+! Copy of 'libgomp.oacc-fortran/mdc-refcount-1-4-1.f90', without 'finalize' clause.
+
+program main
+  use openacc
+  implicit none
+  integer, parameter :: n = 512
+  type mytype
+    integer, allocatable :: a(:)
+  end type mytype
+  type(mytype) :: var
+
+  allocate(var%a(1:n))
+
+  !$acc data create(var)
+
+  call acc_create(var%a)
+  ! After mapping via runtime API call, separately trigger attach action; see <https://github.com/OpenACC/openacc-spec/issues/301>.
+  !$acc enter data attach(var%a)
+
+  if (.not. acc_is_present(var%a)) stop 1
+  if (.not. acc_is_present(var)) stop 2
+
+  !$acc exit data detach(var%a)
+  print *, "CheCKpOInT1"
+  ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
+  !$acc exit data delete(var%a)
+  !TODO { dg-output "(\n|\r\n|\r)libgomp: attach count underflow(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+  !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
+  !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
+  print *, "CheCKpOInT2"
+  ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
+  if (acc_is_present(var%a)) stop 3
+  if (.not. acc_is_present(var)) stop 4
+
+  !$acc end data
+  if (acc_is_present(var%a)) stop 5
+  if (acc_is_present(var)) stop 6
+
+  deallocate(var%a)
+
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/openacc_version-1.f b/libgomp/testsuite/libgomp.oacc-fortran/openacc_version-1.f
index 537212e..36e9844 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/openacc_version-1.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/openacc_version-1.f
@@ -4,6 +4,6 @@
       implicit none
       include "openacc_lib.h"
 
-      if (openacc_version .ne. 201306) STOP 1
+      if (openacc_version .ne. 201711) STOP 1
 
       end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/openacc_version-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/openacc_version-2.f90
index 54f301b..e815bc1 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/openacc_version-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/openacc_version-2.f90
@@ -4,6 +4,6 @@ program main
   use openacc
   implicit none
 
-  if (openacc_version .ne. 201306) STOP 1
+  if (openacc_version .ne. 201711) STOP 1
 
 end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-10.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-10.f90
new file mode 100644
index 0000000..90cca7c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-10.f90
@@ -0,0 +1,52 @@
+! { dg-do run }
+!
+module m
+  implicit none
+contains
+  pure subroutine add_ps_routine(a, b, c)
+    implicit none
+    !$acc routine seq
+    integer, intent(in)  :: a, b
+    integer, intent(out) :: c
+    integer, parameter :: n = 10
+    integer :: i
+
+    do i = 1, n
+       if (i .eq. 5) then
+          c = a + b
+       end if
+    end do
+  end subroutine add_ps_routine
+
+  elemental impure function add_ef(a, b) result(c)
+    implicit none
+    !$acc routine
+    integer, intent(in)  :: a, b
+    integer :: c
+
+    call add_ps_routine(a, b, c)
+  end function add_ef
+end module m
+
+program main
+  use m
+  implicit none
+  integer, parameter :: n = 10
+  integer, dimension(n) :: a_a
+  integer, dimension(n) :: b_a
+  integer, dimension(n) :: c_a
+  integer :: i
+
+  a_a = [(3 * i, i = 1, n)]
+  b_a = [(-2 * i, i = 1, n)]
+  !$acc parallel copyin(a_a, b_a) copyout(c_a)
+  !$acc loop gang
+  do i = 1, n
+     if (i .eq. 4) then
+        c_a = add_ef(a_a, b_a)
+     end if
+  end do
+  !$acc end parallel
+  if (any (c_a /= [(i, i=1, 10)])) stop 1
+  !print *, a
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/stop-1.f b/libgomp/testsuite/libgomp.oacc-fortran/stop-1.f
index af267fc..2c00d2e 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/stop-1.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/stop-1.f
@@ -3,6 +3,10 @@
       PROGRAM MAIN
       IMPLICIT NONE
 
+! Initialize before the checkpoint, in case this produces any output.
+!$ACC PARALLEL
+!$ACC END PARALLEL
+
       PRINT *, "CheCKpOInT"
 !$ACC PARALLEL
       STOP
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/stop-2.f b/libgomp/testsuite/libgomp.oacc-fortran/stop-2.f
index 13c0684..adade54 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/stop-2.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/stop-2.f
@@ -3,6 +3,10 @@
       PROGRAM MAIN
       IMPLICIT NONE
 
+! Initialize before the checkpoint, in case this produces any output.
+!$ACC PARALLEL
+!$ACC END PARALLEL
+
       PRINT *, "CheCKpOInT"
 !$ACC PARALLEL
       STOP 35
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/stop-3.f b/libgomp/testsuite/libgomp.oacc-fortran/stop-3.f
index 3bd7446..157e369 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/stop-3.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/stop-3.f
@@ -3,6 +3,10 @@
       PROGRAM MAIN
       IMPLICIT NONE
 
+! Initialize before the checkpoint, in case this produces any output.
+!$ACC PARALLEL
+!$ACC END PARALLEL
+
       PRINT *, "CheCKpOInT"
 !$ACC PARALLEL
       STOP "SiGN"