Merge branch 'master' into devel/sphinx

23 files changed, 2114 insertions, 122 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index d61a806..b39144e 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,74 @@
+2022-11-04  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* libgomp-plugin.h (OFFLOAD_TARGET_TYPE_INTEL_MIC): Remove.
+	* libgomp.texi (OpenMP Context Selectors): Remove Intel MIC
+	documentation.
+	* plugin/configfrag.ac <enable_offload_targets>
+	[*-intelmic-* | *-intelmicemul-*]: Remove.
+	* configure: Regenerate.
+	* testsuite/lib/libgomp.exp (libgomp_init): Remove 'liboffloadmic'
+	handling.
+	(offload_target_to_openacc_device_type)
+	[$offload_target = *-intelmic*]: Remove.
+	(check_effective_target_offload_device_intel_mic)
+	(check_effective_target_offload_device_any_intel_mic): Remove.
+	* testsuite/libgomp.c-c++-common/on_device_arch.h
+	(device_arch_intel_mic, on_device_arch_intel_mic, any_device_arch)
+	(any_device_arch_intel_mic): Remove.
+	* testsuite/libgomp.c-c++-common/target-45.c: Remove
+	'offload_device_any_intel_mic' XFAIL.
+	* testsuite/libgomp.fortran/target10.f90: Likewise.
+
+2022-11-03  Tobias Burnus  <tobias@codesourcery.com>
+
+	* testsuite/libgomp.fortran/target-11.f90: New test.
+	* testsuite/libgomp.fortran/target-13.f90: New test.
+
+2022-11-02  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR libgomp/106643
+	PR fortran/96668
+	* oacc-mem.c (goacc_enter_data_internal): Support
+	OpenACC 'declare create' with Fortran allocatable arrays, part II.
+	* testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90:
+	Adjust.
+	* testsuite/libgomp.oacc-fortran/pr106643-1.f90: New.
+
+2022-11-02  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR libgomp/106643
+	* oacc-mem.c (goacc_enter_data_internal): Support
+	OpenACC 'declare create' with Fortran allocatable arrays, part I.
+	* testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90:
+	New.
+	* testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90:
+	New.
+
+2022-11-02  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90:
+	New.
+
+2022-11-02  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90:
+	New.
+
+2022-11-02  Cesar Philippidis  <cesar@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+
+	* testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90: New.
+
+2022-10-28  Julian Brown  <julian@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR middle-end/90115
+	* testsuite/libgomp.oacc-fortran/declare-1.f90: Adjust scan output.
+	* testsuite/libgomp.oacc-fortran/host_data-5.F90: Likewise.
+	* testsuite/libgomp.oacc-fortran/if-1.f90: Likewise.
+	* testsuite/libgomp.oacc-fortran/print-1.f90: Likewise.
+	* testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Likewise.
+
 2022-10-24  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* plugin/plugin-nvptx.c (nvptx_open_device): Initialize
diff --git a/libgomp/configure b/libgomp/configure
index 35424d2..45a769e 100755
--- a/libgomp/configure
+++ b/libgomp/configure
@@ -15200,9 +15200,6 @@ if test x"$enable_offload_targets" != x; then
     tgt=`echo $tgt | sed 's/=.*//'`
     tgt_plugin=
     case $tgt in
-      *-intelmic-* | *-intelmicemul-*)
-	tgt_plugin=intelmic
-	;;
       nvptx*)
 	case "${target}" in
 	  aarch64*-*-* | powerpc64le-*-* | x86_64-*-*)
diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h
index 875f967..ac38782 100644
--- a/libgomp/libgomp-plugin.h
+++ b/libgomp/libgomp-plugin.h
@@ -49,7 +49,6 @@ enum offload_target_type
   OFFLOAD_TARGET_TYPE_HOST = 2,
   /* OFFLOAD_TARGET_TYPE_HOST_NONSHM = 3 removed.  */
   OFFLOAD_TARGET_TYPE_NVIDIA_PTX = 5,
-  OFFLOAD_TARGET_TYPE_INTEL_MIC = 6,
   OFFLOAD_TARGET_TYPE_HSA = 7,
   OFFLOAD_TARGET_TYPE_GCN = 8
 };
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 4ff1fdb..3e8340b 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -4303,7 +4303,7 @@ offloading devices (it's not clear if they should be):
 
 @multitable @columnfractions .60 .10 .25
 @headitem @code{arch} @tab @code{kind} @tab @code{isa}
-@item @code{intel_mic}, @code{x86}, @code{x86_64}, @code{i386}, @code{i486},
+@item @code{x86}, @code{x86_64}, @code{i386}, @code{i486},
       @code{i586}, @code{i686}, @code{ia32}
       @tab @code{host}
       @tab See @code{-m...} flags in ``x86 Options'' (without @code{-m})
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 73b2710..233fe0e 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1150,8 +1150,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 	}
       else if (n && groupnum > 1)
 	{
-	  assert (n->refcount != REFCOUNT_INFINITY
-		  && n->refcount != REFCOUNT_LINK);
+	  assert (n->refcount != REFCOUNT_LINK);
 
 	  for (size_t j = i + 1; j <= group_last; j++)
 	    if ((kinds[j] & 0xff) == GOMP_MAP_ATTACH)
@@ -1166,6 +1165,44 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 	  bool processed = false;
 
 	  struct target_mem_desc *tgt = n->tgt;
+
+	  /* Minimal OpenACC variant corresponding to PR96668
+	     "[OpenMP] Re-mapping allocated but previously unallocated
+	     allocatable does not work" 'libgomp/target.c' changes, so that
+	     OpenACC 'declare' code à la PR106643
+	     "[gfortran + OpenACC] Allocate in module causes refcount error"
+	     has a chance to work.  */
+	  if ((kinds[i] & 0xff) == GOMP_MAP_TO_PSET
+	      && tgt->list_count == 0)
+	    {
+	      /* 'declare target'.  */
+	      assert (n->refcount == REFCOUNT_INFINITY);
+
+	      for (size_t k = 1; k < groupnum; k++)
+		{
+		  /* The only thing we expect to see here.  */
+		  assert ((kinds[i + k] & 0xff) == GOMP_MAP_POINTER);
+		}
+
+	      /* Let 'goacc_map_vars' -> 'gomp_map_vars_internal' handle
+		 this.  */
+	      gomp_mutex_unlock (&acc_dev->lock);
+	      struct target_mem_desc *tgt_
+		= goacc_map_vars (acc_dev, aq, groupnum, &hostaddrs[i], NULL,
+				  &sizes[i], &kinds[i], true,
+				  GOMP_MAP_VARS_ENTER_DATA);
+	      assert (tgt_ == NULL);
+	      gomp_mutex_lock (&acc_dev->lock);
+
+	      /* Given that 'goacc_exit_data_internal'/'goacc_exit_datum_1'
+		 will always see 'n->refcount == REFCOUNT_INFINITY',
+		 there's no need to adjust 'n->dynamic_refcount' here.  */
+
+	      processed = true;
+	    }
+	  else
+	    assert (n->refcount != REFCOUNT_INFINITY);
+
 	  for (size_t j = 0; j < tgt->list_count; j++)
 	    if (tgt->list[j].key == n)
 	      {
diff --git a/libgomp/plugin/configfrag.ac b/libgomp/plugin/configfrag.ac
index ab03f94..d3b2589 100644
--- a/libgomp/plugin/configfrag.ac
+++ b/libgomp/plugin/configfrag.ac
@@ -59,9 +59,6 @@ if test x"$enable_offload_targets" != x; then
     tgt=`echo $tgt | sed 's/=.*//'`
     tgt_plugin=
     case $tgt in
-      *-intelmic-* | *-intelmicemul-*)
-	tgt_plugin=intelmic
-	;;
       nvptx*)
 	case "${target}" in
 	  aarch64*-*-* | powerpc64le-*-* | x86_64-*-*)
diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index 4b8c64d..1801fdc 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -119,18 +119,6 @@ proc libgomp_init { args } {
     # Compute what needs to be put into LD_LIBRARY_PATH
     set always_ld_library_path ".:${blddir}/.libs"
 
-    # Add liboffloadmic build directory in LD_LIBRARY_PATH to support
-    # Intel MIC offloading testing.
-    global offload_plugins
-    if { [string match "*,intelmic,*" ",$offload_plugins,"] } {
-	append always_ld_library_path ":${blddir}/../liboffloadmic/.libs"
-	append always_ld_library_path ":${blddir}/../liboffloadmic/plugin/.libs"
-	# libstdc++ is required by liboffloadmic
-	append always_ld_library_path ":${blddir}/../libstdc++-v3/src/.libs"
-	# libgcc_s is required by libstdc++
-	append always_ld_library_path ":${blddir}/../libgcc"
-    }
-
     global offload_additional_lib_paths
     if { $offload_additional_lib_paths != "" } {
 	append always_ld_library_path "${offload_additional_lib_paths}"
@@ -313,9 +301,6 @@ proc offload_target_to_openacc_device_type { offload_target } {
 	disable {
 	    return "host"
 	}
-	*-intelmic* {
-	    return ""
-	}
 	nvptx* {
 	    return "nvidia"
 	}
@@ -449,28 +434,6 @@ proc check_effective_target_openacc_nvidia_accel_selected { } {
     return [string match "nvidia" $openacc_device_type]
 }
 
-# Return 1 if using Intel MIC offload device.
-proc check_effective_target_offload_device_intel_mic { } {
-    return [check_runtime_nocache offload_device_intel_mic {
-      #include "testsuite/libgomp.c-c++-common/on_device_arch.h"
-      int main ()
-	{
-	  return !on_device_arch_intel_mic ();
-	}
-    } ]
-}
-
-# Return 1 if any Intel MIC offload device is available.
-proc check_effective_target_offload_device_any_intel_mic { } {
-    return [check_runtime_nocache offload_device_any_intel_mic {
-      #include "testsuite/libgomp.c-c++-common/on_device_arch.h"
-      int main ()
-	{
-	  return !any_device_arch_intel_mic ();
-	}
-    } ]
-}
-
 # Return 1 if the OpenACC 'host' device type is selected.
 
 proc check_effective_target_openacc_host_selected { } {
diff --git a/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h b/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h
index 6f66dbd..3fb5021 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h
+++ b/libgomp/testsuite/libgomp.c-c++-common/on_device_arch.h
@@ -13,15 +13,8 @@ device_arch_gcn (void)
   return GOMP_DEVICE_GCN;
 }
 
-/* static */ int
-device_arch_intel_mic (void)
-{
-  return GOMP_DEVICE_INTEL_MIC;
-}
-
 #pragma omp declare variant (device_arch_nvptx) match(construct={target},device={arch(nvptx)})
 #pragma omp declare variant (device_arch_gcn) match(construct={target},device={arch(gcn)})
-#pragma omp declare variant (device_arch_intel_mic) match(construct={target},device={arch(intel_mic)})
 /* static */ int
 device_arch (void)
 {
@@ -49,31 +42,3 @@ on_device_arch_gcn ()
 {
   return on_device_arch (GOMP_DEVICE_GCN);
 }
-
-int
-on_device_arch_intel_mic ()
-{
-  return on_device_arch (GOMP_DEVICE_INTEL_MIC);
-}
-
-static int
-any_device_arch (int d)
-{
-  int nd = omp_get_num_devices ();
-  for (int i = 0; i < nd; ++i)
-    {
-      int d_cur;
-      #pragma omp target device(i) map(from:d_cur)
-      d_cur = device_arch ();
-      if (d_cur == d)
-	return 1;
-    }
-
-  return 0;
-}
-
-int
-any_device_arch_intel_mic ()
-{
-  return any_device_arch (GOMP_DEVICE_INTEL_MIC);
-}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-45.c b/libgomp/testsuite/libgomp.c-c++-common/target-45.c
index 27bbedd..73c105d 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/target-45.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-45.c
@@ -1,5 +1,3 @@
-/* { dg-xfail-run-if TODO { offload_device_any_intel_mic } } */
-
 #include <omp.h>
 #include <stdlib.h>
 
diff --git a/libgomp/testsuite/libgomp.fortran/target-11.f90 b/libgomp/testsuite/libgomp.fortran/target-11.f90
new file mode 100644
index 0000000..b0faa2e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-11.f90
@@ -0,0 +1,75 @@
+! Based on libgomp.c/target-23.c
+
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-final { scan-tree-dump "omp target update to\\(xxs\\\[3\\\] \\\[len: 2\\\]\\)" "original" } }
+! { dg-final { scan-tree-dump "omp target update to\\(s\\.s \\\[len: 4\\\]\\)" "original" } }
+! { dg-final { scan-tree-dump "omp target update from\\(s\\.s \\\[len: 4\\\]\\)" "original" } }
+
+module m
+  implicit none
+  type S_type
+    integer s
+    integer, pointer :: u(:) => null()
+    integer :: v(0:4)
+  end type S_type
+  integer, volatile :: z
+end module m
+
+program main
+  use m
+  implicit none
+  integer, target :: u(0:9) = [0, 1, 2, 3, 4, 5, 6, 7, 8, 9]
+  logical :: err
+  type (S_type) :: s
+  integer, pointer :: v(:)
+  integer(kind=2) :: xxs(5)
+  err = .false.
+  s = S_type(9, v=[10, 11, 12, 13, 14])
+  s%u(0:) => u(3:)
+  v(-4+3:) => u(3:)
+  xxs = [-1,-2,-3,-4,-5]
+  !$omp target enter data map (to: s%s, s%u, s%u(0:5)) map (alloc: s%v(1:4), xxs(3:5))
+  s%s = s%s + 1
+  u(3) = u(3) + 1
+  s%v(1) = s%v(1) + 1
+  xxs(3) = -33
+  xxs(4) = -44
+  xxs(5) = -55
+  !$omp target update to (xxs(4))
+  !$omp target update to (s%s) to (s%u(0:2), s%v(1:4))
+
+  !$omp target map (alloc: s%s, s%v(1:4)) map (from: err)
+    err = .false.
+    if (s%s /= 10 .or. s%v(1) /= 12 .or. s%v(2) /= 12 .or. s%v(3) /= 13) &
+      err = .true.
+    if (v(-1) /= 4 .or. v(0) /= 4 .or. v(1) /= 5 .or. v(2) /= 6 .or. v(3) /= 7) &
+      err = .true.
+    if (xxs(4) /= -44) &
+      err = .true.
+    s%s = s%s + 1
+    s%v(2) = s%v(2) + 2
+    v(-1) = 5
+    v(3) = 9
+  !$omp end target
+
+  if (err) &
+    error stop
+
+  !$omp target map (alloc: s%u(0:5))
+    err = .false.
+    if (s%u(0) /= 5 .or. s%u(1) /= 4 .or. s%u(2) /= 5 .or. s%u(3) /= 6 .or. s%u(4) /= 9) &
+      err = .true.
+    s%u(1) = 12
+  !$omp end target
+
+  !$omp target update from (s%s, s%u(0:5)) from (s%v(1:4))
+  if (err .or. s%s /= 11 .or. u(0) /= 0 .or. u(1) /= 1 .or. u(2) /= 2 .or. u(3) /= 5 &
+      .or. u(4) /= 12 .or. u(5) /= 5 .or. u(6) /= 6 .or. u(7) /= 9 .or. u(8) /= 8    &
+      .or. u(9) /= 9 .or. s%v(0) /= 10 .or. s%v(1) /= 12 .or. s%v(2) /= 14           &
+      .or. s%v(3) /= 13 .or. s%v(4) /= 14)                                           &
+    error stop
+  ! !$omp target exit data map (release: s%s)
+  ! !$omp target exit data map (release: s%u(0:5))
+  ! !$omp target exit data map (delete: s%v(1:4))
+  ! !$omp target exit data map (release: s%s)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target-13.f90 b/libgomp/testsuite/libgomp.fortran/target-13.f90
new file mode 100644
index 0000000..6aacc77
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-13.f90
@@ -0,0 +1,159 @@
+module m
+  implicit none
+  type t
+    integer :: s, a(5)
+  end type t
+
+  type t2
+    integer :: s, a(5)
+    type(t) :: st, at(2:3)
+  end type t2
+
+  interface operator(/=)
+    procedure ne_compare_t
+    procedure ne_compare_t2
+  end interface
+
+contains
+
+  logical pure elemental function ne_compare_t (a, b) result(res)
+    type(t), intent(in) :: a, b
+    res = (a%s /= b%s) .or. any(a%a /= b%a)
+  end function
+
+  logical pure elemental function ne_compare_t2 (a, b) result(res)
+    type(t2), intent(in) :: a, b
+    res = (a%s /= b%s) .or. any(a%a /= b%a)     &
+          .or. (a%st /= b%st) .or. any(a%at /= b%at)
+  end function
+end module m
+
+program p
+use m
+implicit none
+
+type(t2) :: var1, var2(5), var3(:)
+type(t2) :: var1a, var2a(5), var3a(:)
+allocatable :: var3, var3a
+logical :: shared_memory = .false.
+
+!$omp target map(to: shared_memory)
+  shared_memory = .true.
+!$omp end target
+
+var1 = T2(1, [1,2,3,4,5], T(11, [11,22,33,44,55]), &
+          [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])])
+
+var2 = [T2(101, [201,202,203,204,205], T(2011, [2011,2022,2033,2044,2055]), &
+           [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]),       &
+        T2(111, [211,212,213,214,215], T(2111, [2111,2122,2133,2144,2155]), &
+           [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]),       &
+        T2(121, [221,222,223,224,225], T(2211, [2211,2222,2233,2244,2255]), &
+           [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]),       &
+        T2(131, [231,232,233,234,235], T(2311, [2311,2322,2333,2344,2355]), &
+           [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]),       &
+        T2(141, [241,242,243,244,245], T(2411, [2411,2422,2433,2444,2455]), &
+           [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])])]
+
+var3 = [T2(301, [401,402,403,404,405], T(4011, [4011,4022,4033,4044,4055]), &
+           [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]),       &
+        T2(311, [411,412,413,414,415], T(4111, [4111,4122,4133,4144,4155]), &
+           [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]),       &
+        T2(321, [421,422,423,424,425], T(4211, [4211,4222,4233,4244,4255]), &
+           [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]),       &
+        T2(331, [431,432,433,434,435], T(4311, [4311,4322,4333,4344,4355]), &
+           [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])]),       &
+        T2(341, [441,442,443,444,445], T(4411, [4411,4422,4433,4444,4455]), &
+           [T(-11, [-11,-22,-33,-44,-55]), T(11, [11,22,33,44,55])])]
+
+var1a = var1
+var2a = var2
+var3a = var3
+
+!$omp target enter data map(to:var1)
+!$omp target enter data map(to:var2)
+!$omp target enter data map(to:var3)
+
+! ---------------
+
+!$omp target update from(var1%at(2:3))
+
+if (var1a /= var1) error stop
+if (any (var2a /= var2)) error stop
+if (any (var3a /= var3)) error stop
+
+! ---------------
+
+!$omp target
+  var1%st%s = 1243
+  var2(3)%at(2) = T(123, [345,64,356,39,13])
+  var2(3)%at(3) = T(48, [74,162,572,357,3])
+!$omp end target
+
+if (.not. shared_memory) then
+  if (var1 /= var1) error stop
+  if (any (var2a /= var2)) error stop
+  if (any (var3a /= var3)) error stop
+endif
+
+!$omp target update from(var1%st) from(var2(3)%at(2:3))
+
+var1a%st%s = 1243
+var2a(3)%at(2) = T(123, [345,64,356,39,13])
+var2a(3)%at(3) = T(48, [74,162,572,357,3])
+if (var1 /= var1) error stop
+if (any (var2a /= var2)) error stop
+if (any (var3a /= var3)) error stop
+
+! ---------------
+
+var3(1) = var2(1)
+var1%at(2)%a = var2(1)%a
+var1%at(3)%a = var2(2)%a
+
+var1a = var1
+var2a = var2
+var3a = var3
+
+!$omp target update to(var3) to(var1%at(2:3))
+
+!$omp target
+  var3(1)%s = var3(1)%s + 123
+  var1%at(2)%a = var1%at(2)%a * 7
+  var1%at(3)%s = var1%at(3)%s * (-3)
+!$omp end target
+
+if (.not. shared_memory) then
+  if (var1 /= var1) error stop
+  if (any (var2a /= var2)) error stop
+  if (any (var3a /= var3)) error stop
+endif
+
+var3a(1)%s = var3a(1)%s + 123
+var1a%at(2)%a = var1a%at(2)%a * 7
+var1a%at(3)%s = var1a%at(3)%s * (-3)
+
+block
+  integer, volatile :: i1,i2,i3,i4
+  i1 = 1
+  i2 = 2
+  i3 = 1
+  i4 = 2
+  !$omp target update from(var3(i1:i2)) from(var1%at(i3:i4))
+  i1 = 3
+  i2 = 3
+  i3 = 1
+  i4 = 5
+  !$omp target update from(var1%at(i1)%s) from(var1%at(i2)%a(i3:i4))
+end block
+
+if (var1 /= var1) error stop
+if (any (var2a /= var2)) error stop
+if (any (var3a /= var3)) error stop
+
+! ---------------
+
+!$omp target exit data map(from:var1)
+!$omp target exit data map(from:var2)
+!$omp target exit data map(from:var3)
+end
diff --git a/libgomp/testsuite/libgomp.fortran/target10.f90 b/libgomp/testsuite/libgomp.fortran/target10.f90
index 3145255..4876cce 100644
--- a/libgomp/testsuite/libgomp.fortran/target10.f90
+++ b/libgomp/testsuite/libgomp.fortran/target10.f90
@@ -1,5 +1,4 @@
 ! { dg-do run }
-! { dg-xfail-run-if TODO { offload_device_any_intel_mic } }
 
 program main
   use omp_lib
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
index 51776a1..89bd4a2 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
@@ -215,7 +215,7 @@ program main
   ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } .-1 }
   ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-2 }
   ! { dg-note {variable 'S\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
-  ! { dg-note {variable 'desc\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-4 }
+  ! { dg-note {variable 'desc\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-4 }
   use vars
   use openacc
   implicit none
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90
new file mode 100644
index 0000000..759873b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90
@@ -0,0 +1,278 @@
+! Test OpenACC 'declare create' with allocatable arrays.
+
+! { dg-do run }
+
+!TODO-OpenACC-declare-allocate
+! Missing support for OpenACC "Changes from Version 2.0 to 2.5":
+! "The 'declare create' directive with a Fortran 'allocatable' has new behavior".
+! Thus, after 'allocate'/before 'deallocate', do
+! '!$acc enter data create'/'!$acc exit data delete' manually.
+
+!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.
+
+! { dg-additional-options -fopt-info-all-omp }
+! { dg-additional-options -foffload=-fopt-info-all-omp }
+
+! { dg-additional-options --param=openacc-privatization=noisy }
+! { dg-additional-options -foffload=--param=openacc-privatization=noisy }
+! Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
+! { dg-prune-output {note: variable '[Di]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
+
+! { dg-additional-options -Wopenacc-parallelism }
+
+! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
+! passed to 'incr' may be unset, and in that case, it will be set to [...]",
+! so to maintain compatibility with earlier Tcl releases, we manually
+! initialize counter variables:
+! { dg-line l_dummy[variable c 0] }
+! { dg-message dummy {} { target iN-VAl-Id } l_dummy } to avoid
+! "WARNING: dg-line var l_dummy defined, but not used".
+
+
+module vars
+  implicit none
+  integer, parameter :: n = 100
+  real*8, allocatable :: b(:)
+ !$acc declare create (b)
+end module vars
+
+program test
+  use vars
+  use openacc
+  implicit none
+  real*8 :: a
+  integer :: i
+
+  interface
+     subroutine sub1
+       !$acc routine gang
+     end subroutine sub1
+
+     subroutine sub2
+     end subroutine sub2
+
+     real*8 function fun1 (ix)
+       integer ix
+       !$acc routine seq
+     end function fun1
+
+     real*8 function fun2 (ix)
+       integer ix
+       !$acc routine seq
+     end function fun2
+  end interface
+
+  if (allocated (b)) error stop
+
+  ! Test local usage of an allocated declared array.
+
+  allocate (b(n))
+  !$acc enter data create (b)
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  a = 2.0
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c }
+  ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = i * a
+  end do
+
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc update host(b)
+
+  do i = 1, n
+     if (b(i) /= i*a) error stop
+  end do
+
+  !$acc exit data delete (b)
+  deallocate (b)
+
+  ! Test the usage of an allocated declared array inside an acc
+  ! routine subroutine.
+
+  allocate (b(n))
+  !$acc enter data create (b)
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc parallel
+  call sub1 ! { dg-line l[incr c] }
+  ! { dg-optimized {assigned OpenACC gang worker vector loop parallelism} {} { target *-*-* } l$c }
+  !$acc end parallel
+
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc update host(b)
+
+  do i = 1, n
+     if (b(i) /= i*2) error stop
+  end do
+
+  !$acc exit data delete (b)
+  deallocate (b)
+
+  ! Test the usage of an allocated declared array inside a host
+  ! subroutine.
+
+  call sub2
+
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc update host(b)
+
+  do i = 1, n
+     if (b(i) /= 1.0) error stop
+  end do
+
+  !$acc exit data delete (b)
+  deallocate (b)
+
+  if (allocated (b)) error stop
+
+  ! Test the usage of an allocated declared array inside an acc
+  ! routine function.
+
+  allocate (b(n))
+  !$acc enter data create (b)
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c }
+  ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = 1.0
+  end do
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c }
+  ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = fun1 (i) ! { dg-line l[incr c] }
+     ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l$c }
+  end do
+
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc update host(b)
+
+  do i = 1, n
+     if (b(i) /= i) error stop
+  end do
+
+  !$acc exit data delete (b)
+  deallocate (b)
+
+  ! Test the usage of an allocated declared array inside a host
+  ! function.
+
+  allocate (b(n))
+  !$acc enter data create (b)
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c }
+  ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = 1.0
+  end do
+
+  !$acc update host(b)
+
+  do i = 1, n
+     b(i) = fun2 (i)
+  end do
+
+  if (.not.acc_is_present (b)) error stop
+
+  do i = 1, n
+     if (b(i) /= i*i) error stop
+  end do
+
+  !$acc exit data delete (b)
+  deallocate (b)
+end program test ! { dg-line l[incr c] }
+! { dg-bogus {note: variable 'overflow\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c }
+! { dg-bogus {note: variable 'not_prev_allocated\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c }
+! { dg-bogus {note: variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {TODO n/a} { xfail *-*-* } l$c }
+
+! Set each element in array 'b' at index i to i*2.
+
+subroutine sub1 ! { dg-line subroutine_sub1 }
+  use vars
+  implicit none
+  integer i
+  !$acc routine gang
+  ! { dg-bogus {[Ww]arning: region is worker partitioned but does not contain worker partitioned code} {TODO default 'gang' 'vector'} { xfail *-*-* } subroutine_sub1 }
+
+  !$acc loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = i*2
+  end do
+end subroutine sub1
+
+! Allocate array 'b', and set it to all 1.0.
+
+subroutine sub2
+  use vars
+  use openacc
+  implicit none
+  integer i
+
+  allocate (b(n))
+  !$acc enter data create (b)
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = 1.0
+  end do
+end subroutine sub2
+
+! Return b(i) * i;
+
+real*8 function fun1 (i)
+  use vars
+  implicit none
+  integer i
+  !$acc routine seq
+
+  fun1 = b(i) * i
+end function fun1
+
+! Return b(i) * i * i;
+
+real*8 function fun2 (i)
+  use vars
+  implicit none
+  integer i
+
+  fun2 = b(i) * i * i
+end function fun2
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90
new file mode 100644
index 0000000..e4cb9c3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90
@@ -0,0 +1,278 @@
+! Test OpenACC 'declare create' with allocatable arrays.
+
+! { dg-do run }
+
+!TODO-OpenACC-declare-allocate
+! Missing support for OpenACC "Changes from Version 2.0 to 2.5":
+! "The 'declare create' directive with a Fortran 'allocatable' has new behavior".
+! Thus, after 'allocate'/before 'deallocate', call 'acc_create'/'acc_delete'
+! manually.
+
+!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.
+
+! { dg-additional-options -fopt-info-all-omp }
+! { dg-additional-options -foffload=-fopt-info-all-omp }
+
+! { dg-additional-options --param=openacc-privatization=noisy }
+! { dg-additional-options -foffload=--param=openacc-privatization=noisy }
+! Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
+! { dg-prune-output {note: variable '[Di]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
+
+! { dg-additional-options -Wopenacc-parallelism }
+
+! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
+! passed to 'incr' may be unset, and in that case, it will be set to [...]",
+! so to maintain compatibility with earlier Tcl releases, we manually
+! initialize counter variables:
+! { dg-line l_dummy[variable c 0] }
+! { dg-message dummy {} { target iN-VAl-Id } l_dummy } to avoid
+! "WARNING: dg-line var l_dummy defined, but not used".
+
+
+module vars
+  implicit none
+  integer, parameter :: n = 100
+  real*8, allocatable :: b(:)
+ !$acc declare create (b)
+end module vars
+
+program test
+  use vars
+  use openacc
+  implicit none
+  real*8 :: a
+  integer :: i
+
+  interface
+     subroutine sub1
+       !$acc routine gang
+     end subroutine sub1
+
+     subroutine sub2
+     end subroutine sub2
+
+     real*8 function fun1 (ix)
+       integer ix
+       !$acc routine seq
+     end function fun1
+
+     real*8 function fun2 (ix)
+       integer ix
+       !$acc routine seq
+     end function fun2
+  end interface
+
+  if (allocated (b)) error stop
+
+  ! Test local usage of an allocated declared array.
+
+  allocate (b(n))
+  call acc_create (b)
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  a = 2.0
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c }
+  ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = i * a
+  end do
+
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc update host(b)
+
+  do i = 1, n
+     if (b(i) /= i*a) error stop
+  end do
+
+  call acc_delete (b)
+  deallocate (b)
+
+  ! Test the usage of an allocated declared array inside an acc
+  ! routine subroutine.
+
+  allocate (b(n))
+  call acc_create (b)
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc parallel
+  call sub1 ! { dg-line l[incr c] }
+  ! { dg-optimized {assigned OpenACC gang worker vector loop parallelism} {} { target *-*-* } l$c }
+  !$acc end parallel
+
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc update host(b)
+
+  do i = 1, n
+     if (b(i) /= i*2) error stop
+  end do
+
+  call acc_delete (b)
+  deallocate (b)
+
+  ! Test the usage of an allocated declared array inside a host
+  ! subroutine.
+
+  call sub2
+
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc update host(b)
+
+  do i = 1, n
+     if (b(i) /= 1.0) error stop
+  end do
+
+  call acc_delete (b)
+  deallocate (b)
+
+  if (allocated (b)) error stop
+
+  ! Test the usage of an allocated declared array inside an acc
+  ! routine function.
+
+  allocate (b(n))
+  call acc_create (b)
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c }
+  ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = 1.0
+  end do
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c }
+  ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = fun1 (i) ! { dg-line l[incr c] }
+     ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l$c }
+  end do
+
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc update host(b)
+
+  do i = 1, n
+     if (b(i) /= i) error stop
+  end do
+
+  call acc_delete (b)
+  deallocate (b)
+
+  ! Test the usage of an allocated declared array inside a host
+  ! function.
+
+  allocate (b(n))
+  call acc_create (b)
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c }
+  ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = 1.0
+  end do
+
+  !$acc update host(b)
+
+  do i = 1, n
+     b(i) = fun2 (i)
+  end do
+
+  if (.not.acc_is_present (b)) error stop
+
+  do i = 1, n
+     if (b(i) /= i*i) error stop
+  end do
+
+  call acc_delete (b)
+  deallocate (b)
+end program test ! { dg-line l[incr c] }
+! { dg-bogus {note: variable 'overflow\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c }
+! { dg-bogus {note: variable 'not_prev_allocated\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c }
+! { dg-bogus {note: variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {TODO n/a} { xfail *-*-* } l$c }
+
+! Set each element in array 'b' at index i to i*2.
+
+subroutine sub1 ! { dg-line subroutine_sub1 }
+  use vars
+  implicit none
+  integer i
+  !$acc routine gang
+  ! { dg-bogus {[Ww]arning: region is worker partitioned but does not contain worker partitioned code} {TODO default 'gang' 'vector'} { xfail *-*-* } subroutine_sub1 }
+
+  !$acc loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = i*2
+  end do
+end subroutine sub1
+
+! Allocate array 'b', and set it to all 1.0.
+
+subroutine sub2
+  use vars
+  use openacc
+  implicit none
+  integer i
+
+  allocate (b(n))
+  call acc_create (b)
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = 1.0
+  end do
+end subroutine sub2
+
+! Return b(i) * i;
+
+real*8 function fun1 (i)
+  use vars
+  implicit none
+  integer i
+  !$acc routine seq
+
+  fun1 = b(i) * i
+end function fun1
+
+! Return b(i) * i * i;
+
+real*8 function fun2 (i)
+  use vars
+  implicit none
+  integer i
+
+  fun2 = b(i) * i * i
+end function fun2
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90
new file mode 100644
index 0000000..1c8ccd9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90
@@ -0,0 +1,268 @@
+! Test OpenACC 'declare create' with allocatable arrays.
+
+! { dg-do run }
+
+!TODO-OpenACC-declare-allocate
+! Not currently implementing correct '-DACC_MEM_SHARED=0' behavior:
+! Missing support for OpenACC "Changes from Version 2.0 to 2.5":
+! "The 'declare create' directive with a Fortran 'allocatable' has new behavior".
+! { dg-xfail-run-if TODO { *-*-* } { -DACC_MEM_SHARED=0 } }
+
+!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.
+
+! { dg-additional-options -fopt-info-all-omp }
+! { dg-additional-options -foffload=-fopt-info-all-omp }
+
+! { dg-additional-options --param=openacc-privatization=noisy }
+! { dg-additional-options -foffload=--param=openacc-privatization=noisy }
+! Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
+! { dg-prune-output {note: variable '[Di]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} }
+
+! { dg-additional-options -Wopenacc-parallelism }
+
+! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
+! passed to 'incr' may be unset, and in that case, it will be set to [...]",
+! so to maintain compatibility with earlier Tcl releases, we manually
+! initialize counter variables:
+! { dg-line l_dummy[variable c 0] }
+! { dg-message dummy {} { target iN-VAl-Id } l_dummy } to avoid
+! "WARNING: dg-line var l_dummy defined, but not used".
+
+
+module vars
+  implicit none
+  integer, parameter :: n = 100
+  real*8, allocatable :: b(:)
+ !$acc declare create (b)
+end module vars
+
+program test
+  use vars
+  use openacc
+  implicit none
+  real*8 :: a
+  integer :: i
+
+  interface
+     subroutine sub1
+       !$acc routine gang
+     end subroutine sub1
+
+     subroutine sub2
+     end subroutine sub2
+
+     real*8 function fun1 (ix)
+       integer ix
+       !$acc routine seq
+     end function fun1
+
+     real*8 function fun2 (ix)
+       integer ix
+       !$acc routine seq
+     end function fun2
+  end interface
+
+  if (allocated (b)) error stop
+
+  ! Test local usage of an allocated declared array.
+
+  allocate (b(n))
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  a = 2.0
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c }
+  ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = i * a
+  end do
+
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc update host(b)
+
+  do i = 1, n
+     if (b(i) /= i*a) error stop
+  end do
+
+  deallocate (b)
+
+  ! Test the usage of an allocated declared array inside an acc
+  ! routine subroutine.
+
+  allocate (b(n))
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc parallel
+  call sub1 ! { dg-line l[incr c] }
+  ! { dg-optimized {assigned OpenACC gang worker vector loop parallelism} {} { target *-*-* } l$c }
+  !$acc end parallel
+
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc update host(b)
+
+  do i = 1, n
+     if (b(i) /= i*2) error stop
+  end do
+
+  deallocate (b)
+
+  ! Test the usage of an allocated declared array inside a host
+  ! subroutine.
+
+  call sub2
+
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc update host(b)
+
+  do i = 1, n
+     if (b(i) /= 1.0) error stop
+  end do
+
+  deallocate (b)
+
+  if (allocated (b)) error stop
+
+  ! Test the usage of an allocated declared array inside an acc
+  ! routine function.
+
+  allocate (b(n))
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c }
+  ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = 1.0
+  end do
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c }
+  ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = fun1 (i) ! { dg-line l[incr c] }
+     ! { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l$c }
+  end do
+
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc update host(b)
+
+  do i = 1, n
+     if (b(i) /= i) error stop
+  end do
+
+  deallocate (b)
+
+  ! Test the usage of an allocated declared array inside a host
+  ! function.
+
+  allocate (b(n))
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' ought to be adjusted for OpenACC privatization level: 'vector'} {} { target *-*-* } l$c }
+  !   { dg-note {variable 'i' adjusted for OpenACC privatization level: 'vector'} {} { target { ! openacc_host_selected } } l$c }
+  ! { dg-note {variable 'i\.[0-9]+' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = 1.0
+  end do
+
+  !$acc update host(b)
+
+  do i = 1, n
+     b(i) = fun2 (i)
+  end do
+
+  if (.not.acc_is_present (b)) error stop
+
+  do i = 1, n
+     if (b(i) /= i*i) error stop
+  end do
+
+  deallocate (b)
+end program test ! { dg-line l[incr c] }
+! { dg-bogus {note: variable 'overflow\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c }
+! { dg-bogus {note: variable 'not_prev_allocated\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {TODO n/a} { xfail *-*-* } l$c }
+! { dg-bogus {note: variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {TODO n/a} { xfail *-*-* } l$c }
+
+! Set each element in array 'b' at index i to i*2.
+
+subroutine sub1 ! { dg-line subroutine_sub1 }
+  use vars
+  implicit none
+  integer i
+  !$acc routine gang
+  ! { dg-bogus {[Ww]arning: region is worker partitioned but does not contain worker partitioned code} {TODO default 'gang' 'vector'} { xfail *-*-* } subroutine_sub1 }
+
+  !$acc loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = i*2
+  end do
+end subroutine sub1
+
+! Allocate array 'b', and set it to all 1.0.
+
+subroutine sub2
+  use vars
+  use openacc
+  implicit none
+  integer i
+
+  allocate (b(n))
+
+  if (.not.allocated (b)) error stop
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc parallel loop ! { dg-line l[incr c] }
+  ! { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l$c }
+  ! { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l$c }
+  do i = 1, n
+     b(i) = 1.0
+  end do
+end subroutine sub2
+
+! Return b(i) * i;
+
+real*8 function fun1 (i)
+  use vars
+  implicit none
+  integer i
+  !$acc routine seq
+
+  fun1 = b(i) * i
+end function fun1
+
+! Return b(i) * i * i;
+
+real*8 function fun2 (i)
+  use vars
+  implicit none
+  integer i
+
+  fun2 = b(i) * i * i
+end function fun2
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90
new file mode 100644
index 0000000..6604f72
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90
@@ -0,0 +1,438 @@
+! Test OpenACC 'declare create' with allocatable arrays.
+
+! { dg-do run }
+
+! Note that we're not testing OpenACC semantics here, but rather documenting
+! current GCC behavior, specifically, behavior concerning updating of
+! host/device array descriptors.
+! { dg-skip-if n/a { *-*-* } { -DACC_MEM_SHARED=1 } }
+
+!TODO-OpenACC-declare-allocate
+! Missing support for OpenACC "Changes from Version 2.0 to 2.5":
+! "The 'declare create' directive with a Fortran 'allocatable' has new behavior".
+! Thus, after 'allocate'/before 'deallocate', do
+! '!$acc enter data create'/'!$acc exit data delete' manually.
+
+
+!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.
+
+
+!TODO OpenACC 'serial' vs. GCC/nvptx:
+!TODO { dg-prune-output {using 'vector_length \(32\)', ignoring 1} }
+
+
+! { dg-additional-options -fdump-tree-original }
+! { dg-additional-options -fdump-tree-gimple }
+
+
+module vars
+  implicit none
+  integer, parameter :: n1_lb = -3
+  integer, parameter :: n1_ub = 6
+  integer, parameter :: n2_lb = -9999
+  integer, parameter :: n2_ub = 22222
+
+  integer, allocatable :: b(:)
+  !$acc declare create (b)
+
+end module vars
+
+program test
+  use vars
+  use openacc
+  implicit none
+  integer :: i
+
+  ! Identifiers for purposes of reliable '-fdump-tree-[...]' scanning.
+  integer :: id1_1, id1_2
+
+  interface
+
+     subroutine verify_initial
+       implicit none
+       !$acc routine seq
+     end subroutine verify_initial
+
+     subroutine verify_n1_allocated
+       implicit none
+       !$acc routine seq
+     end subroutine verify_n1_allocated
+
+     subroutine verify_n1_values (addend)
+       implicit none
+       !$acc routine gang
+       integer, value :: addend
+     end subroutine verify_n1_values
+
+     subroutine verify_n1_deallocated (expect_allocated)
+       implicit none
+       !$acc routine seq
+       logical, value :: expect_allocated
+     end subroutine verify_n1_deallocated
+
+     subroutine verify_n2_allocated
+       implicit none
+       !$acc routine seq
+     end subroutine verify_n2_allocated
+
+     subroutine verify_n2_values (addend)
+       implicit none
+       !$acc routine gang
+       integer, value :: addend
+     end subroutine verify_n2_values
+
+     subroutine verify_n2_deallocated (expect_allocated)
+       implicit none
+       !$acc routine seq
+       logical, value :: expect_allocated
+     end subroutine verify_n2_deallocated
+
+  end interface
+
+  call acc_create (id1_1)
+  call acc_create (id1_2)
+
+  call verify_initial
+  ! It is important here (and similarly, following) that there is no data
+  ! clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'.
+  !$acc serial
+  call verify_initial
+  !$acc end serial
+
+  allocate (b(n1_lb:n1_ub))
+  call verify_n1_allocated
+  if (acc_is_present (b)) error stop
+  !$acc enter data create (b)
+  ! This is now OpenACC "present":
+  if (.not.acc_is_present (b)) error stop
+  ! ..., and got the actual array descriptor installed:
+  !$acc serial
+  call verify_n1_allocated
+  !$acc end serial
+
+  do i = n1_lb, n1_ub
+     b(i) = i - 1
+  end do
+
+  ! In 'declare-allocatable-array_descriptor-1-runtime.f90', this does "verify
+  ! that host-to-device copy doesn't touch the device-side (still initial)
+  ! array descriptor (but it does copy the array data").  This is here not
+  ! applicable anymore, as we've already gotten the actual array descriptor
+  ! installed.  Thus now verify that it does copy the array data.
+  call acc_update_device (b)
+  !$acc serial
+  call verify_n1_allocated
+  !$acc end serial
+
+  b = 40
+
+  !$acc parallel copyout (id1_1) ! No data clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'.
+  call verify_n1_values (-1)
+  id1_1 = 0
+  !$acc end parallel
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(from:id1_1\)$} 1 original } }
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(from:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
+
+  !$acc parallel copy (b) copyout (id1_2)
+  ! As already present, 'copy (b)' doesn't copy; addend is still '-1'.
+  call verify_n1_values (-1)
+  id1_2 = 0
+  !$acc end parallel
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(tofrom:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2\)$} 1 original } }
+  !TODO ..., but without an actual use of 'b', the gimplifier removes the
+  !TODO 'GOMP_MAP_TO_PSET':
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(tofrom:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+
+  ! In 'declare-allocatable-array_descriptor-1-runtime.f90', this does "verify
+  ! that device-to-host copy doesn't touch the host-side array descriptor,
+  ! doesn't copy out the device-side (still initial) array descriptor (but it
+  ! does copy the array data)".  This is here not applicable anymore, as we've
+  ! already gotten the actual array descriptor installed.  Thus now verify that
+  ! it does copy the array data.
+  call acc_update_self (b)
+  call verify_n1_allocated
+
+  do i = n1_lb, n1_ub
+     if (b(i) /= i - 1) error stop
+     b(i) = b(i) + 2
+  end do
+
+  ! The same using the OpenACC 'update' directive.
+
+  !$acc update device (b) self (id1_1)
+  ! We do have 'GOMP_MAP_TO_PSET' here:
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc update map\(force_to:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_1\);$} 1 original } }
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_update map\(force_to:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
+  ! ..., but it's silently skipped in 'GOACC_update'.
+  !$acc serial
+  call verify_n1_allocated
+  !$acc end serial
+
+  b = 41
+
+  !$acc parallel
+  call verify_n1_values (1)
+  !$acc end parallel
+
+  !$acc parallel copy (b)
+  call verify_n1_values (1)
+  !$acc end parallel
+
+  !$acc update self (b) self (id1_2)
+  ! We do have 'GOMP_MAP_TO_PSET' here:
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc update map\(force_from:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_2\);$} 1 original } }
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_update map\(force_from:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+  ! ..., but it's silently skipped in 'GOACC_update'.
+  call verify_n1_allocated
+
+  do i = n1_lb, n1_ub
+     if (b(i) /= i + 1) error stop
+     b(i) = b(i) + 2
+  end do
+
+  ! Now test that (potentially re-)installing the actual array descriptor is a
+  ! no-op, via a data clause for 'b' (explicit or implicit): must get a
+  ! 'GOMP_MAP_TO_PSET'.
+  !$acc serial present (b) copyin (id1_2)
+  call verify_n1_allocated
+  !TODO Use of 'b':
+  id1_2 = ubound (b, 1)
+  !$acc end serial
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2\)$} 1 original } }
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+
+  !$acc parallel copyin (id1_1) ! No data clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'.
+  call verify_n1_values (1)
+  id1_1 = 0
+  !$acc end parallel
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(to:id1_1\)$} 1 original } }
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(to:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
+
+  !$acc parallel copy (b) copyin (id1_2)
+  ! As already present, 'copy (b)' doesn't copy; addend is still '1'.
+  call verify_n1_values (1)
+  id1_2 = 0
+  !$acc end parallel
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(tofrom:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2\)$} 1 original } }
+  !TODO ..., but without an actual use of 'b', the gimplifier removes the
+  !TODO 'GOMP_MAP_TO_PSET':
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(tofrom:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+
+  call verify_n1_allocated
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc exit data delete (b)
+  if (.not.allocated (b)) error stop
+  if (acc_is_present (b)) error stop
+  ! The device-side array descriptor doesn't get updated, so 'b' still appears
+  ! as "allocated":
+  !$acc serial
+  call verify_n1_allocated
+  !$acc end serial
+
+  deallocate (b)
+  call verify_n1_deallocated (.false.)
+  ! The device-side array descriptor doesn't get updated, so 'b' still appears
+  ! as "allocated":
+  !$acc serial
+  call verify_n1_allocated
+  !$acc end serial
+
+  ! Now try to install the actual array descriptor, via a data clause for 'b'
+  ! (explicit or implicit): must get a 'GOMP_MAP_TO_PSET', which then in
+  ! 'gomp_map_vars_internal' is handled as 'declare target', but because of
+  ! '*(void **) hostaddrs[i] == NULL', we've got 'has_always_ptrset == false',
+  ! 'always_to_cnt == 0', and therefore 'gomp_map_vars_existing' doesn't update
+  ! the 'GOMP_MAP_TO_PSET'.
+  ! The device-side array descriptor doesn't get updated, so 'b' still appears
+  ! as "allocated":
+  !TODO Why does 'present (b)' still work here?
+  !$acc serial present (b) copyout (id1_2)
+  call verify_n1_deallocated (.true.)
+  !TODO Use of 'b'.
+  id1_2 = ubound (b, 1)
+  !$acc end serial
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2\)$} 1 original } }
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+
+
+  ! Restart the procedure, with different array dimensions.
+
+  allocate (b(n2_lb:n2_ub))
+  call verify_n2_allocated
+  if (acc_is_present (b)) error stop
+  !$acc enter data create (b)
+  if (.not.acc_is_present (b)) error stop
+  ! ..., and got the actual array descriptor installed:
+  !$acc serial
+  call verify_n2_allocated
+  !$acc end serial
+
+  do i = n2_lb, n2_ub
+     b(i) = i + 20
+  end do
+
+  call acc_update_device (b)
+  !$acc serial
+  call verify_n2_allocated
+  !$acc end serial
+
+  b = -40
+
+  !$acc parallel
+  call verify_n2_values (20)
+  !$acc end parallel
+
+  !$acc parallel copy (b)
+  call verify_n2_values (20)
+  !$acc end parallel
+
+  call acc_update_self (b)
+  call verify_n2_allocated
+
+  do i = n2_lb, n2_ub
+     if (b(i) /= i + 20) error stop
+     b(i) = b(i) - 40
+  end do
+
+  !$acc update device (b)
+  !$acc serial
+  call verify_n2_allocated
+  !$acc end serial
+
+  b = -41
+
+  !$acc parallel
+  call verify_n2_values (-20)
+  !$acc end parallel
+
+  !$acc parallel copy (b)
+  call verify_n2_values (-20)
+  !$acc end parallel
+
+  !$acc update self (b)
+  call verify_n2_allocated
+
+  do i = n2_lb, n2_ub
+     if (b(i) /= i - 20) error stop
+     b(i) = b(i) + 10
+  end do
+
+  !$acc serial present (b) copy (id1_2)
+  call verify_n2_allocated
+  !TODO Use of 'b':
+  id1_2 = ubound (b, 1)
+  !$acc end serial
+
+  !$acc parallel
+  call verify_n2_values (-20)
+  !$acc end parallel
+
+  !$acc parallel copy (b)
+  call verify_n2_values (-20)
+  !$acc end parallel
+
+  call verify_n2_allocated
+  if (.not.acc_is_present (b)) error stop
+
+  !$acc exit data delete (b)
+  if (.not.allocated (b)) error stop
+  if (acc_is_present (b)) error stop
+  !$acc serial
+  call verify_n2_allocated
+  !$acc end serial
+
+  deallocate (b)
+  call verify_n2_deallocated (.false.)
+  !$acc serial
+  call verify_n2_allocated
+  !$acc end serial
+
+  !$acc serial present (b) copy (id1_2)
+  call verify_n2_deallocated (.true.)
+  !TODO Use of 'b':
+  id1_2 = ubound (b, 1)
+  !$acc end serial
+
+end program test
+
+
+subroutine verify_initial
+  use vars
+  implicit none
+  !$acc routine seq
+
+  if (allocated (b)) error stop "verify_initial allocated"
+  if (any (lbound (b) /= [0])) error stop "verify_initial lbound"
+  if (any (ubound (b) /= [0])) error stop "verify_initial ubound"
+end subroutine verify_initial
+
+subroutine verify_n1_allocated
+  use vars
+  implicit none
+  !$acc routine seq
+
+  if (.not.allocated (b)) error stop "verify_n1_allocated allocated"
+  if (any (lbound (b) /= [n1_lb])) error stop "verify_n1_allocated lbound"
+  if (any (ubound (b) /= [n1_ub])) error stop "verify_n1_allocated ubound"
+end subroutine verify_n1_allocated
+
+subroutine verify_n1_values (addend)
+  use vars
+  implicit none
+  !$acc routine gang
+  integer, value :: addend
+  integer :: i
+
+  !$acc loop
+  do i = n1_lb, n1_ub
+     if (b(i) /= i + addend) error stop
+  end do
+end subroutine verify_n1_values
+
+subroutine verify_n1_deallocated (expect_allocated)
+  use vars
+  implicit none
+  !$acc routine seq
+  logical, value :: expect_allocated
+
+  if (allocated(b) .neqv. expect_allocated) error stop "verify_n1_deallocated allocated"
+  ! Apparently 'deallocate'ing doesn't unset the bounds.
+  if (any (lbound (b) /= [n1_lb])) error stop "verify_n1_deallocated lbound"
+  if (any (ubound (b) /= [n1_ub])) error stop "verify_n1_deallocated ubound"
+end subroutine verify_n1_deallocated
+
+subroutine verify_n2_allocated
+  use vars
+  implicit none
+  !$acc routine seq
+
+  if (.not.allocated(b)) error stop "verify_n2_allocated allocated"
+  if (any (lbound (b) /= [n2_lb])) error stop "verify_n2_allocated lbound"
+  if (any (ubound (b) /= [n2_ub])) error stop "verify_n2_allocated ubound"
+end subroutine verify_n2_allocated
+
+subroutine verify_n2_values (addend)
+  use vars
+  implicit none
+  !$acc routine gang
+  integer, value :: addend
+  integer :: i
+
+  !$acc loop
+  do i = n2_lb, n2_ub
+     if (b(i) /= i + addend) error stop
+  end do
+end subroutine verify_n2_values
+
+subroutine verify_n2_deallocated (expect_allocated)
+  use vars
+  implicit none
+  !$acc routine seq
+  logical, value :: expect_allocated
+
+  if (allocated(b) .neqv. expect_allocated) error stop "verify_n2_deallocated allocated"
+  ! Apparently 'deallocate'ing doesn't unset the bounds.
+  if (any (lbound (b) /= [n2_lb])) error stop "verify_n2_deallocated lbound"
+  if (any (ubound (b) /= [n2_ub])) error stop "verify_n2_deallocated ubound"
+end subroutine verify_n2_deallocated
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90
new file mode 100644
index 0000000..b27f312
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90
@@ -0,0 +1,402 @@
+! Test OpenACC 'declare create' with allocatable arrays.
+
+! { dg-do run }
+
+! Note that we're not testing OpenACC semantics here, but rather documenting
+! current GCC behavior, specifically, behavior concerning updating of
+! host/device array descriptors.
+! { dg-skip-if n/a { *-*-* } { -DACC_MEM_SHARED=1 } }
+
+!TODO-OpenACC-declare-allocate
+! Missing support for OpenACC "Changes from Version 2.0 to 2.5":
+! "The 'declare create' directive with a Fortran 'allocatable' has new behavior".
+! Thus, after 'allocate'/before 'deallocate', call 'acc_create'/'acc_delete'
+! manually.
+
+
+!TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.
+
+
+!TODO OpenACC 'serial' vs. GCC/nvptx:
+!TODO { dg-prune-output {using 'vector_length \(32\)', ignoring 1} }
+
+
+! { dg-additional-options -fdump-tree-original }
+! { dg-additional-options -fdump-tree-gimple }
+
+
+module vars
+  implicit none
+  integer, parameter :: n1_lb = -3
+  integer, parameter :: n1_ub = 6
+  integer, parameter :: n2_lb = -9999
+  integer, parameter :: n2_ub = 22222
+
+  integer, allocatable :: b(:)
+  !$acc declare create (b)
+
+end module vars
+
+program test
+  use vars
+  use openacc
+  implicit none
+  integer :: i
+
+  ! Identifiers for purposes of reliable '-fdump-tree-[...]' scanning.
+  integer :: id1_1, id1_2
+
+  interface
+
+     subroutine verify_initial
+       implicit none
+       !$acc routine seq
+     end subroutine verify_initial
+
+     subroutine verify_n1_allocated
+       implicit none
+       !$acc routine seq
+     end subroutine verify_n1_allocated
+
+     subroutine verify_n1_values (addend)
+       implicit none
+       !$acc routine gang
+       integer, value :: addend
+     end subroutine verify_n1_values
+
+     subroutine verify_n1_deallocated (expect_allocated)
+       implicit none
+       !$acc routine seq
+       logical, value :: expect_allocated
+     end subroutine verify_n1_deallocated
+
+     subroutine verify_n2_allocated
+       implicit none
+       !$acc routine seq
+     end subroutine verify_n2_allocated
+
+     subroutine verify_n2_values (addend)
+       implicit none
+       !$acc routine gang
+       integer, value :: addend
+     end subroutine verify_n2_values
+
+     subroutine verify_n2_deallocated (expect_allocated)
+       implicit none
+       !$acc routine seq
+       logical, value :: expect_allocated
+     end subroutine verify_n2_deallocated
+
+  end interface
+
+  call acc_create (id1_1)
+  call acc_create (id1_2)
+
+  call verify_initial
+  ! It is important here (and similarly, following) that there is no data
+  ! clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'.
+  !$acc serial
+  call verify_initial
+  !$acc end serial
+
+  allocate (b(n1_lb:n1_ub))
+  call verify_n1_allocated
+  if (acc_is_present (b)) error stop
+  call acc_create (b)
+  ! This is now OpenACC "present":
+  if (.not.acc_is_present (b)) error stop
+  ! This still has the initial array descriptor:
+  !$acc serial
+  call verify_initial
+  !$acc end serial
+
+  do i = n1_lb, n1_ub
+     b(i) = i - 1
+  end do
+
+  ! Verify that host-to-device copy doesn't touch the device-side (still
+  ! initial) array descriptor (but it does copy the array data).
+  call acc_update_device (b)
+  !$acc serial
+  call verify_initial
+  !$acc end serial
+
+  b = 40
+
+  ! Verify that device-to-host copy doesn't touch the host-side array
+  ! descriptor, doesn't copy out the device-side (still initial) array
+  ! descriptor (but it does copy the array data).
+  call acc_update_self (b)
+  call verify_n1_allocated
+
+  do i = n1_lb, n1_ub
+     if (b(i) /= i - 1) error stop
+     b(i) = b(i) + 2
+  end do
+
+  ! The same using the OpenACC 'update' directive.
+
+  !$acc update device (b) self (id1_1)
+  ! We do have 'GOMP_MAP_TO_PSET' here:
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc update map\(force_to:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_1\);$} 1 original } }
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_update map\(force_to:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
+  ! ..., but it's silently skipped in 'GOACC_update'.
+  !$acc serial
+  call verify_initial
+  !$acc end serial
+
+  b = 41
+
+  !$acc update self (b) self (id1_2)
+  ! We do have 'GOMP_MAP_TO_PSET' here:
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc update map\(force_from:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_2\);$} 1 original } }
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_update map\(force_from:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+  ! ..., but it's silently skipped in 'GOACC_update'.
+  call verify_n1_allocated
+
+  do i = n1_lb, n1_ub
+     if (b(i) /= i + 1) error stop
+     b(i) = b(i) + 2
+  end do
+
+  ! Now install the actual array descriptor, via a data clause for 'b'
+  ! (explicit or implicit): must get a 'GOMP_MAP_TO_PSET', which then in
+  ! 'gomp_map_vars_internal' is handled as 'declare target', and because of
+  ! '*(void **) hostaddrs[i] != NULL', we've got 'has_always_ptrset == true',
+  ! 'always_to_cnt == 1', and therefore 'gomp_map_vars_existing' does update
+  ! the 'GOMP_MAP_TO_PSET'.
+  !$acc serial present (b) copyin (id1_1)
+  call verify_initial
+  id1_1 = 0
+  !$acc end serial
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_1\)$} 1 original } }
+  !TODO ..., but without an actual use of 'b', the gimplifier removes the
+  !TODO 'GOMP_MAP_TO_PSET':
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
+  !$acc serial present (b) copyin (id1_2)
+  call verify_n1_allocated
+  !TODO Use of 'b':
+  id1_2 = ubound (b, 1)
+  !$acc end serial
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2\)$} 1 original } }
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+
+  !$acc parallel copyin (id1_1) ! No data clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'.
+  call verify_n1_values (1)
+  id1_1 = 0
+  !$acc end parallel
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(to:id1_1\)$} 1 original } }
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(to:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
+
+  !$acc parallel copy (b) copyin (id1_2)
+  ! As already present, 'copy (b)' doesn't copy; addend is still '1'.
+  call verify_n1_values (1)
+  id1_2 = 0
+  !$acc end parallel
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(tofrom:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2\)$} 1 original } }
+  !TODO ..., but without an actual use of 'b', the gimplifier removes the
+  !TODO 'GOMP_MAP_TO_PSET':
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(tofrom:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+
+  call verify_n1_allocated
+  if (.not.acc_is_present (b)) error stop
+
+  call acc_delete (b)
+  if (.not.allocated (b)) error stop
+  if (acc_is_present (b)) error stop
+  ! The device-side array descriptor doesn't get updated, so 'b' still appears
+  ! as "allocated":
+  !$acc serial
+  call verify_n1_allocated
+  !$acc end serial
+
+  deallocate (b)
+  call verify_n1_deallocated (.false.)
+  ! The device-side array descriptor doesn't get updated, so 'b' still appears
+  ! as "allocated":
+  !$acc serial
+  call verify_n1_allocated
+  !$acc end serial
+
+  ! Now try to install the actual array descriptor, via a data clause for 'b'
+  ! (explicit or implicit): must get a 'GOMP_MAP_TO_PSET', which then in
+  ! 'gomp_map_vars_internal' is handled as 'declare target', but because of
+  ! '*(void **) hostaddrs[i] == NULL', we've got 'has_always_ptrset == false',
+  ! 'always_to_cnt == 0', and therefore 'gomp_map_vars_existing' doesn't update
+  ! the 'GOMP_MAP_TO_PSET'.
+  ! The device-side array descriptor doesn't get updated, so 'b' still appears
+  ! as "allocated":
+  !TODO Why does 'present (b)' still work here?
+  !$acc serial present (b) copyout (id1_2)
+  call verify_n1_deallocated (.true.)
+  !TODO Use of 'b'.
+  id1_2 = ubound (b, 1)
+  !$acc end serial
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2\)$} 1 original } }
+  ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
+
+
+  ! Restart the procedure, with different array dimensions.
+
+  allocate (b(n2_lb:n2_ub))
+  call verify_n2_allocated
+  if (acc_is_present (b)) error stop
+  call acc_create (b)
+  if (.not.acc_is_present (b)) error stop
+  ! This still has the previous (n1) array descriptor:
+  !$acc serial
+  call verify_n1_deallocated (.true.)
+  !$acc end serial
+
+  do i = n2_lb, n2_ub
+     b(i) = i + 20
+  end do
+
+  call acc_update_device (b)
+  !$acc serial
+  call verify_n1_deallocated (.true.)
+  !$acc end serial
+
+  b = -40
+
+  call acc_update_self (b)
+  call verify_n2_allocated
+
+  do i = n2_lb, n2_ub
+     if (b(i) /= i + 20) error stop
+     b(i) = b(i) - 40
+  end do
+
+  !$acc update device (b)
+  !$acc serial
+  call verify_n1_deallocated (.true.)
+  !$acc end serial
+
+  b = -41
+
+  !$acc update self (b)
+  call verify_n2_allocated
+
+  do i = n2_lb, n2_ub
+     if (b(i) /= i - 20) error stop
+     b(i) = b(i) + 10
+  end do
+
+  !$acc serial present (b) copy (id1_2)
+  call verify_n2_allocated
+  !TODO Use of 'b':
+  id1_2 = ubound (b, 1)
+  !$acc end serial
+
+  !$acc parallel
+  call verify_n2_values (-20)
+  !$acc end parallel
+
+  !$acc parallel copy (b)
+  call verify_n2_values (-20)
+  !$acc end parallel
+
+  call verify_n2_allocated
+  if (.not.acc_is_present (b)) error stop
+
+  call acc_delete (b)
+  if (.not.allocated (b)) error stop
+  if (acc_is_present (b)) error stop
+  !$acc serial
+  call verify_n2_allocated
+  !$acc end serial
+
+  deallocate (b)
+  call verify_n2_deallocated (.false.)
+  !$acc serial
+  call verify_n2_allocated
+  !$acc end serial
+
+  !$acc serial present (b) copy (id1_2)
+  call verify_n2_deallocated (.true.)
+  !TODO Use of 'b':
+  id1_2 = ubound (b, 1)
+  !$acc end serial
+
+end program test
+
+
+subroutine verify_initial
+  use vars
+  implicit none
+  !$acc routine seq
+
+  if (allocated (b)) error stop "verify_initial allocated"
+  if (any (lbound (b) /= [0])) error stop "verify_initial lbound"
+  if (any (ubound (b) /= [0])) error stop "verify_initial ubound"
+end subroutine verify_initial
+
+subroutine verify_n1_allocated
+  use vars
+  implicit none
+  !$acc routine seq
+
+  if (.not.allocated (b)) error stop "verify_n1_allocated allocated"
+  if (any (lbound (b) /= [n1_lb])) error stop "verify_n1_allocated lbound"
+  if (any (ubound (b) /= [n1_ub])) error stop "verify_n1_allocated ubound"
+end subroutine verify_n1_allocated
+
+subroutine verify_n1_values (addend)
+  use vars
+  implicit none
+  !$acc routine gang
+  integer, value :: addend
+  integer :: i
+
+  !$acc loop
+  do i = n1_lb, n1_ub
+     if (b(i) /= i + addend) error stop
+  end do
+end subroutine verify_n1_values
+
+subroutine verify_n1_deallocated (expect_allocated)
+  use vars
+  implicit none
+  !$acc routine seq
+  logical, value :: expect_allocated
+
+  if (allocated(b) .neqv. expect_allocated) error stop "verify_n1_deallocated allocated"
+  ! Apparently 'deallocate'ing doesn't unset the bounds.
+  if (any (lbound (b) /= [n1_lb])) error stop "verify_n1_deallocated lbound"
+  if (any (ubound (b) /= [n1_ub])) error stop "verify_n1_deallocated ubound"
+end subroutine verify_n1_deallocated
+
+subroutine verify_n2_allocated
+  use vars
+  implicit none
+  !$acc routine seq
+
+  if (.not.allocated(b)) error stop "verify_n2_allocated allocated"
+  if (any (lbound (b) /= [n2_lb])) error stop "verify_n2_allocated lbound"
+  if (any (ubound (b) /= [n2_ub])) error stop "verify_n2_allocated ubound"
+end subroutine verify_n2_allocated
+
+subroutine verify_n2_values (addend)
+  use vars
+  implicit none
+  !$acc routine gang
+  integer, value :: addend
+  integer :: i
+
+  !$acc loop
+  do i = n2_lb, n2_ub
+     if (b(i) /= i + addend) error stop
+  end do
+end subroutine verify_n2_values
+
+subroutine verify_n2_deallocated (expect_allocated)
+  use vars
+  implicit none
+  !$acc routine seq
+  logical, value :: expect_allocated
+
+  if (allocated(b) .neqv. expect_allocated) error stop "verify_n2_deallocated allocated"
+  ! Apparently 'deallocate'ing doesn't unset the bounds.
+  if (any (lbound (b) /= [n2_lb])) error stop "verify_n2_deallocated lbound"
+  if (any (ubound (b) /= [n2_ub])) error stop "verify_n2_deallocated ubound"
+end subroutine verify_n2_deallocated
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90 b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
index 93e9ee0..c3453a5 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
@@ -43,7 +43,7 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
   ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target { ! openacc_host_selected } } .-2 }
   ! { dg-note {variable 'p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
   ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target { ! openacc_host_selected } } .-5 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-5 }
 #if !ACC_MEM_SHARED
     if (acc_is_present(p, c_sizeof(p))) stop 5
     if (acc_is_present(parr, 1)) stop 6
@@ -54,8 +54,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
     ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
     ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
     ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
-    ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-6 }
-    ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
+    ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-6 }
+    ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-8 }
       ! not mapped yet, so it will be equal to the host pointer.
       if (transfer(c_loc(p), host_p) /= host_p) stop 7
@@ -74,9 +74,9 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
     ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
     ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 }
     ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } .-7 }
-    ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 }
-    ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-9 }
-    ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-10 }
+    ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 }
+    ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-9 }
+    ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-10 }
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-11 }
       if (.not. acc_is_present(p, c_sizeof(p))) stop 11
       if (.not. acc_is_present(parr, 1)) stop 12
@@ -90,8 +90,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
       ! { dg-note {variable 'host_p\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-3 }
       ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
       ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
-      ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-6 }
-      ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
+      ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-6 }
+      ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
       ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-8 }
 #if ACC_MEM_SHARED
         if (transfer(c_loc(p), host_p) /= host_p) stop 15
@@ -110,8 +110,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
         ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
         ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
         ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 }
-        ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
-        ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 }
+        ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
+        ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 }
 #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
@@ -129,8 +129,8 @@ subroutine foo (p2, parr, host_p, host_parr, cond)
         ! { dg-note {variable 'parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-4 }
         ! { dg-note {variable 'host_parr\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-5 }
         ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-6 }
-        ! { dg-note {variable 'D\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-7 }
-        ! { dg-note {variable 'transfer\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "TODO" { target *-*-* } .-8 }
+        ! { dg-note {variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-7 }
+        ! { dg-note {variable 'transfer\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } .-8 }
 #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
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
index c6d6764..e0cfd91 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90
@@ -382,7 +382,7 @@ program main
   b(:) = 1.0
 
   !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1)
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-1 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 }
 
 #if !ACC_MEM_SHARED
   if (acc_is_present (a) .eqv. .TRUE.) STOP 21
@@ -396,7 +396,7 @@ program main
 
   !$acc data copyin (a(1:N)) if (1 == 1)
   ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 
 #if !ACC_MEM_SHARED
     if (acc_is_present (a) .eqv. .FALSE.) STOP 23
@@ -404,7 +404,7 @@ program main
 
     !$acc data copyout (b(1:N)) if (0 == 1)
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-    ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+    ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 #if !ACC_MEM_SHARED
       if (acc_is_present (b) .eqv. .TRUE.) STOP 24
 #endif
@@ -877,7 +877,7 @@ program main
   b(:) = 1.0
 
   !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1)
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-1 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 }
 
 #if !ACC_MEM_SHARED
   if (acc_is_present (a) .eqv. .TRUE.) STOP 56
@@ -891,7 +891,7 @@ program main
 
   !$acc data copyin (a(1:N)) if (1 == 1)
   ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-  ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+  ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 
 #if !ACC_MEM_SHARED
     if (acc_is_present (a) .eqv. .FALSE.) STOP 58
@@ -899,7 +899,7 @@ program main
 
     !$acc data copyout (b(1:N)) if (0 == 1)
     ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 }
-    ! { dg-note {variable 'parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target { ! openacc_host_selected } } .-2 }
+    ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 }
 #if !ACC_MEM_SHARED
       if (acc_is_present (b) .eqv. .TRUE.) STOP 59
 #endif
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr106643-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr106643-1.f90
new file mode 100644
index 0000000..a9c969e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr106643-1.f90
@@ -0,0 +1,83 @@
+! { dg-do run }
+! { dg-additional-options -cpp }
+
+
+!TODO OpenACC 'serial' vs. GCC/nvptx:
+!TODO { dg-prune-output {using 'vector_length \(32\)', ignoring 1} }
+
+
+module m_macron
+
+    implicit none
+
+    real(kind(0d0)), allocatable, dimension(:) :: valls
+    !$acc declare create(valls)
+
+contains
+
+    subroutine s_macron_compute(size)
+
+        integer :: size
+
+        !$acc routine seq
+
+#if ACC_MEM_SHARED
+        if (valls(size) /= 1) error stop
+#else
+        if (valls(size) /= size - 2) error stop
+#endif
+
+        valls(size) = size + 2
+
+    end subroutine s_macron_compute
+
+    subroutine s_macron_init(size)
+
+        integer :: size
+
+        print*, "size=", size
+
+        print*, "allocate(valls(1:size))"
+        allocate(valls(1:size))
+
+        print*, "acc enter data create(valls(1:size))"
+        !$acc enter data create(valls(1:size))
+
+        print*, "!$acc update device(valls(1:size))"
+        valls(size) = size - 2
+        !$acc update device(valls(1:size))
+
+        valls(size) = 1
+
+        !$acc serial
+        call s_macron_compute(size)
+        !$acc end serial
+
+        valls(size) = -1
+
+        !$acc update host(valls(1:size))
+#if ACC_MEM_SHARED
+        if (valls(size) /= -1) error stop
+#else
+        if (valls(size) /= size + 2) error stop
+#endif
+
+        print*, valls(1:size)
+
+        print*, "acc exit data delete(valls)"
+        !$acc exit data delete(valls)
+
+    end subroutine s_macron_init
+
+end module m_macron
+
+
+program p_main
+
+    use m_macron
+
+    implicit none
+
+    call s_macron_init(10)
+
+end program p_main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
index 42a8538..d2f89d9 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90
@@ -6,15 +6,6 @@
 ! Separate file 'print-1-nvptx.f90' for nvptx offloading.
 ! { dg-skip-if "separate file" { offload_target_nvptx } }
 
-! For GCN offloading compilation, when gang-privatizing 'dt_parm.N'
-! (see below), we run into an 'gang-private data-share memory exhausted'
-! error: the default '-mgang-private-size' is too small.  Per
-! 'gcc/fortran/trans-io.cc'/'libgfortran/io/io.h', that one is
-! 'struct st_parameter_dt', which indeed is rather big.  Instead of
-! working out its exact size (which may vary per GCC configuration),
-! raise '-mgang-private-size' to an arbitrary high value.
-! { dg-additional-options "-foffload-options=amdgcn-amdhsa=-mgang-private-size=13579" { target openacc_radeon_accel_selected } }
-
 ! { dg-additional-options "-fopt-info-note-omp" }
 ! { dg-additional-options "-foffload=-fopt-info-note-omp" }
 
@@ -36,9 +27,7 @@ program main
   integer :: var = 42
 
 !$acc parallel ! { dg-line l_compute[incr c_compute] }
-  ! { dg-note {variable 'dt_parm\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute }
-  !   { dg-note {variable 'dt_parm\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} {} { target *-*-* } l_compute$c_compute }
-  !   { dg-note {variable 'dt_parm\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} {} { target { ! openacc_host_selected } } l_compute$c_compute }
+  ! { dg-note {variable 'dt_parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} {} { target *-*-* } l_compute$c_compute }
   write (0, '("The answer is ", I2)') var
 !$acc end parallel
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
index b31f406..498ef70 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90
@@ -122,9 +122,7 @@ contains
     ! { dg-note {variable 'str' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'str' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'str' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
+    ! { dg-note {variable 'char\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
     ! { dg-message {sorry, unimplemented: target cannot support alloca} PR65181 { target openacc_nvidia_accel_selected } l_loop$c_loop }
     do i = 1, 10
       str(i:i) = achar(ichar('A') + i)
@@ -167,9 +165,7 @@ contains
     ! { dg-note {variable 'scalar' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'scalar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
     ! { dg-note {variable 'scalar' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
-    ! { dg-note {variable 'char\.[0-9]+' adjusted for OpenACC privatization level: 'gang'} "" { target { ! { openacc_host_selected || { openacc_nvidia_accel_selected && __OPTIMIZE__ } } } } l_loop$c_loop }
+    ! { dg-note {variable 'char\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target *-*-* } l_loop$c_loop }
     do i = 1, 15
       scalar(i:i) = achar(ichar('A') + i)
     end do