Merge branch 'master' into devel/sphinx

14 files changed, 579 insertions, 16 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index ea62dc6..d3b4758 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,47 @@
+2022-09-14  Julian Brown  <julian@codesourcery.com>
+
+	* testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c: New test.
+	* testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c: New test.
+	* testsuite/libgomp.oacc-c++/deep-copy-17.C: New test.
+	* testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c: Move
+	test to here, make "run" test.
+
+2022-09-13  Jakub Jelinek  <jakub@redhat.com>
+
+	PR libgomp/106906
+	* env.c (get_icv_member_addr): Cast false to void * before assigning
+	it to icv_addr[1], and comment the whole assignment out.
+
+2022-09-13  Tobias Burnus  <tobias@codesourcery.com>
+
+	* libgomp.texi (gcn): Move misplaced -march=sm_30 remark to ...
+	(nvptx): ... here.
+
+2022-09-12  Tobias Burnus  <tobias@codesourcery.com>
+
+	* libgomp.texi (Offload-Target Specifics: nvptx): Document
+	that reverse offload requires >= -march=sm_35.
+	* testsuite/libgomp.c-c++-common/requires-4.c: Build for nvptx
+	with -misa=sm_35.
+	* testsuite/libgomp.c-c++-common/requires-5.c: Likewise.
+	* testsuite/libgomp.c-c++-common/requires-6.c: Likewise.
+	* testsuite/libgomp.c-c++-common/reverse-offload-1.c: Likewise.
+	* testsuite/libgomp.fortran/reverse-offload-1.f90: Likewise.
+	* testsuite/libgomp.c/reverse-offload-sm30.c: New test.
+
+2022-09-12  Tobias Burnus  <tobias@codesourcery.com>
+
+	* libgomp.texi (OpenMP 5.1 Impl. Status): Add two new minor items.
+	(OpenMP 5.2 Impl. Status): Improve omp/omx/ompx wording.
+
+2022-09-12  Jakub Jelinek  <jakub@redhat.com>
+
+	PR libgomp/106894
+	* testsuite/libgomp.c-c++-common/icv-6.c: Include string.h.
+	(main): Avoid tests for which corresponding non-_ALL suffixed variable
+	is in the environment, or for OMP_NUM_TEAMS on the device
+	OMP_NUM_TEAMS_DEV_?.
+
 2022-09-10  Iain Sandoe  <iain@sandoe.co.uk>
 
 	* env.c (initialize_env): Include libiberty environ.h.
diff --git a/libgomp/env.c b/libgomp/env.c
index ac8c764..0249966 100644
--- a/libgomp/env.c
+++ b/libgomp/env.c
@@ -1892,14 +1892,14 @@ get_icv_member_addr (struct gomp_initial_icvs *icvs, int icv_code,
     {
     case GOMP_ICV_NTEAMS:
       icv_addr[0] = &icvs->nteams_var;
-      icv_addr[1] = false;
+      /* icv_addr[1] = (void *) false; */
       break;
     case GOMP_ICV_DYNAMIC:
       icv_addr[0] = &(*icvs).dyn_var;
       break;
     case GOMP_ICV_TEAMS_THREAD_LIMIT:
       icv_addr[0] = &icvs->teams_thread_limit_var;
-      icv_addr[1] = false;
+      /* icv_addr[1] = (void *) false; */
       break;
     case GOMP_ICV_SCHEDULE:
       icv_addr[0] = &icvs->run_sched_var;
@@ -1907,7 +1907,7 @@ get_icv_member_addr (struct gomp_initial_icvs *icvs, int icv_code,
       break;
     case GOMP_ICV_THREAD_LIMIT:
       icv_addr[0] = &icvs->thread_limit_var;
-      icv_addr[1] = false;
+      /* icv_addr[1] = (void *) false; */
       icv_addr[2] = (void *) UINT_MAX;
       break;
     case GOMP_ICV_NTHREADS:
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index ce3ba76..addf2d4 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -348,6 +348,9 @@ The OpenMP 4.5 specification is fully supported.
 @item Support @code{begin/end declare target} syntax in C/C++ @tab N @tab
 @item Pointer predetermined firstprivate getting initialized
 to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab
+@item @code{begin declare target} directive @tab N @tab
+@item For Fortran, diagnose placing declarative before/between @code{USE},
+      @code{IMPORT}, and @code{IMPLICIT} as invalid @tab N @tab
 @end multitable
 
 
@@ -362,12 +365,13 @@ to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab
       @tab N @tab
 @item @code{omp}/@code{ompx}/@code{omx} sentinels and @code{omp_}/@code{ompx_}
       namespaces @tab N/A
-      @tab warning for @code{omp/ompx} sentinels@footnote{@code{omp/ompx}
-      sentinels as C/C++ pragma and C++ attributes are warned for with
+      @tab warning for @code{ompx/omx} sentinels@footnote{The @code{ompx}
+      sentinel as C/C++ pragma and C++ attributes are warned for with
       @code{-Wunknown-pragmas} (implied by @code{-Wall}) and @code{-Wattributes}
       (enabled by default), respectively; for Fortran free-source code, there is
-      a warning enabled by default and for fixed-source code with
-      @code{-Wsurprising} (enabled by @code{-Wall})}
+      a warning enabled by default and, for fixed-source code, the @code{omx}
+      sentinel is warned for with with @code{-Wsurprising} (enabled by
+      @code{-Wall}).  Unknown clauses are always rejected with an error.}
 @item Clauses on @code{end} directive can be on directive @tab N @tab
 @item Deprecation of no-argument @code{destroy} clause on @code{depobj}
       @tab N @tab
@@ -4428,6 +4432,9 @@ The implementation remark:
 @item I/O within OpenMP target regions and OpenACC parallel/kernels is supported
       using the C library @code{printf} functions. Note that the Fortran
       @code{print}/@code{write} statements are not supported, yet.
+@item Compilation OpenMP code that contains @code{requires reverse_offload}
+      requires at least @code{-march=sm_35}, compiling for @code{-march=sm_30}
+      is not supported.
 @end itemize
 
 
diff --git a/libgomp/testsuite/libgomp.c-c++-common/icv-6.c b/libgomp/testsuite/libgomp.c-c++-common/icv-6.c
index 7151bd1..e199a18 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/icv-6.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/icv-6.c
@@ -17,6 +17,7 @@
 
 #include <omp.h>
 #include <stdlib.h>
+#include <string.h>
 
 int
 main ()
@@ -25,21 +26,28 @@ main ()
   int chunk_size;
   omp_get_schedule(&kind, &chunk_size);
 
-  if (omp_get_max_teams () != 42
-      || !omp_get_dynamic ()
-      || kind != 3 || chunk_size != 4
-      || omp_get_teams_thread_limit () != 44
-      || omp_get_thread_limit () != 45
-      || omp_get_max_threads () != 46
-      || omp_get_proc_bind () != omp_proc_bind_spread
-      || omp_get_max_active_levels () != 47)
+  if ((!getenv ("OMP_NUM_TEAMS") && omp_get_max_teams () != 42)
+      || (!getenv ("OMP_DYNAMIC") && !omp_get_dynamic ())
+      || (!getenv ("OMP_SCHEDULE") && (kind != 3 || chunk_size != 4))
+      || (!getenv ("OMP_TEAMS_THREAD_LIMIT") && omp_get_teams_thread_limit () != 44)
+      || (!getenv ("OMP_THREAD_LIMIT") && omp_get_thread_limit () != 45)
+      || (!getenv ("OMP_NUM_THREADS") && omp_get_max_threads () != 46)
+      || (!getenv ("OMP_PROC_BIND") && omp_get_proc_bind () != omp_proc_bind_spread)
+      || (!getenv ("OMP_MAX_ACTIVE_LEVELS") && omp_get_max_active_levels () != 47))
     abort ();
 
   int num_devices = omp_get_num_devices () > 3 ? 3 : omp_get_num_devices ();
-  for (int i=0; i < num_devices; i++)
+  for (int i = 0; i < num_devices; i++)
+    {
+      char name[sizeof ("OMP_NUM_TEAMS_DEV_1")];
+      strcpy (name, "OMP_NUM_TEAMS_DEV_1");
+      name[sizeof ("OMP_NUM_TEAMS_DEV_1") - 2] = '0' + i;
+      if (getenv (name))
+	continue;
     #pragma omp target device (i)
       if (omp_get_max_teams () != 43)
 	abort ();
+    }
 
   return 0;
 }
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
index 6ed5a5f..5883eff 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-4.c
@@ -1,4 +1,5 @@
 /* { dg-additional-options "-flto" } */
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
 /* { dg-additional-sources requires-4-aux.c } */
 
 /* Check no diagnostic by device-compiler's or host compiler's lto1.
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
index 7fe0c73..d43d78d 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-5.c
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
 /* { dg-additional-sources requires-5-aux.c } */
 
 /* Depending on offload device capabilities, it may print something like the
diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-6.c b/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
index b00c745..a25b4d2 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-6.c
@@ -1,3 +1,5 @@
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
+
 #pragma omp requires unified_shared_memory, unified_address, reverse_offload
 
 /* The requires line is not active as there is none of:
diff --git a/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c b/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c
index 976e129..52d828c 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/reverse-offload-1.c
@@ -1,4 +1,5 @@
 /* { dg-do run }  */
+/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
 /* { dg-additional-sources reverse-offload-1-aux.c } */
 
 /* Check that reverse offload works in particular:
diff --git a/libgomp/testsuite/libgomp.c/reverse-offload-sm30.c b/libgomp/testsuite/libgomp.c/reverse-offload-sm30.c
new file mode 100644
index 0000000..fbfeae1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/reverse-offload-sm30.c
@@ -0,0 +1,15 @@
+/* { dg-do link { target { offload_target_nvptx } } } */
+/* { dg-additional-options "-foffload-options=nvptx-none=-march=sm_30 -foffload=-mptx=_" } */
+
+#pragma omp requires reverse_offload
+
+int
+main ()
+{
+  #pragma omp target
+    {
+    }
+  return 0;
+}
+
+/* { dg-warning "'omp requires reverse_offload' requires at least 'sm_35' for '-foffload-options=nvptx-none=-march=' - disabling offload-code generation for this device type" "" { target *-*-* } 0 } */
diff --git a/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90 b/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90
index 7cfb8b6..de68011 100644
--- a/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/reverse-offload-1.f90
@@ -1,4 +1,5 @@
 ! { dg-do run }
+! { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } }
 ! { dg-additional-sources reverse-offload-1-aux.f90 }
 
 ! Check that reverse offload works in particular:
diff --git a/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C
new file mode 100644
index 0000000..dacbb52
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C
@@ -0,0 +1,101 @@
+#include <cassert>
+
+/* Test attach/detach operation with pointers and references to structs.  */
+
+typedef struct mystruct {
+  int *a;
+  int b;
+  int *c;
+  int d;
+  int *e;
+} mystruct;
+
+void str (void)
+{
+  int a[10], c[10], e[10];
+  mystruct m = { .a = a, .c = c, .e = e };
+  a[0] = 5;
+  c[0] = 7;
+  e[0] = 9;
+  #pragma acc parallel copy(m.a[0:10], m.b, m.c[0:10], m.d, m.e[0:10])
+  {
+    m.a[0] = m.c[0] + m.e[0];
+  }
+  assert (m.a[0] == 7 + 9);
+}
+
+void strp (void)
+{
+  int *a = new int[10];
+  int *c = new int[10];
+  int *e = new int[10];
+  mystruct *m = new mystruct;
+  m->a = a;
+  m->c = c;
+  m->e = e;
+  a[0] = 6;
+  c[0] = 8;
+  e[0] = 10;
+  #pragma acc parallel copy(m->a[0:10], m->b, m->c[0:10], m->d, m->e[0:10])
+  {
+    m->a[0] = m->c[0] + m->e[0];
+  }
+  assert (m->a[0] == 8 + 10);
+  delete m;
+  delete[] a;
+  delete[] c;
+  delete[] e;
+}
+
+void strr (void)
+{
+  int *a = new int[10];
+  int *c = new int[10];
+  int *e = new int[10];
+  mystruct m;
+  mystruct &n = m;
+  n.a = a;
+  n.c = c;
+  n.e = e;
+  a[0] = 7;
+  c[0] = 9;
+  e[0] = 11;
+  #pragma acc parallel copy(n.a[0:10], n.b, n.c[0:10], n.d, n.e[0:10])
+  {
+    n.a[0] = n.c[0] + n.e[0];
+  }
+  assert (n.a[0] == 9 + 11);
+  delete[] a;
+  delete[] c;
+  delete[] e;
+}
+
+void strrp (void)
+{
+  int a[10], c[10], e[10];
+  mystruct *m = new mystruct;
+  mystruct *&n = m;
+  n->a = a;
+  n->b = 3;
+  n->c = c;
+  n->d = 5;
+  n->e = e;
+  a[0] = 8;
+  c[0] = 10;
+  e[0] = 12;
+  #pragma acc parallel copy(n->a[0:10], n->c[0:10], n->e[0:10])
+  {
+    n->a[0] = n->c[0] + n->e[0];
+  }
+  assert (n->a[0] == 10 + 12);
+  delete m;
+}
+
+int main (int argc, char *argv[])
+{
+  str ();
+  strp ();
+  strr ();
+  strrp ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c
new file mode 100644
index 0000000..27fe1a9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c
@@ -0,0 +1,68 @@
+#include <stdlib.h>
+
+/* Test multiple struct dereferences on one directive, and slices starting at
+   non-zero.  */
+
+typedef struct {
+  int *a;
+  int *b;
+  int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct *m = (mystruct *) malloc (sizeof (*m));
+  int i;
+
+  m->a = (int *) malloc (N * sizeof (int));
+  m->b = (int *) malloc (N * sizeof (int));
+  m->c = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      m->a[i] = 0;
+      m->b[i] = 0;
+      m->c[i] = 0;
+    }
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+#pragma acc parallel loop copy(m->a[0:N])
+      for (j = 0; j < N; j++)
+	m->a[j]++;
+#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10])
+      for (j = 0; j < N; j++)
+	{
+	  m->b[j]++;
+	  if (j > 5 && j < N - 5)
+	    m->c[j]++;
+	}
+    }
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 99)
+	abort ();
+      if (m->b[i] != 99)
+	abort ();
+      if (i > 5 && i < N-5)
+	{
+	  if (m->c[i] != 99)
+	    abort ();
+	}
+      else
+	{
+	  if (m->c[i] != 0)
+	    abort ();
+	}
+    }
+
+  free (m->a);
+  free (m->b);
+  free (m->c);
+  free (m);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c
new file mode 100644
index 0000000..a7308e8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c
@@ -0,0 +1,231 @@
+#include <stdlib.h>
+
+/* Test mapping chained indirect struct accesses, mixed in different ways.  */
+
+typedef struct {
+  int *a;
+  int b;
+  int *c;
+} str1;
+
+typedef struct {
+  int d;
+  int *e;
+  str1 *f;
+} str2;
+
+typedef struct {
+  int g;
+  int h;
+  str2 *s2;
+} str3;
+
+typedef struct {
+  str3 m;
+  str3 n;
+} str4;
+
+void
+zero_arrays (str4 *s, int N)
+{
+  for (int i = 0; i < N; i++)
+    {
+      s->m.s2->e[i] = 0;
+      s->m.s2->f->a[i] = 0;
+      s->m.s2->f->c[i] = 0;
+      s->n.s2->e[i] = 0;
+      s->n.s2->f->a[i] = 0;
+      s->n.s2->f->c[i] = 0;
+    }
+}
+
+void
+alloc_s2 (str2 **s, int N)
+{
+  (*s) = (str2 *) malloc (sizeof (str2));
+  (*s)->f = (str1 *) malloc (sizeof (str1));
+  (*s)->e = (int *) malloc (sizeof (int) * N);
+  (*s)->f->a = (int *) malloc (sizeof (int) * N);
+  (*s)->f->c = (int *) malloc (sizeof (int) * N);
+}
+
+int main (int argc, char* argv[])
+{
+  const int N = 1024;
+  str4 p, *q;
+  int i;
+
+  alloc_s2 (&p.m.s2, N);
+  alloc_s2 (&p.n.s2, N);
+  q = (str4 *) malloc (sizeof (str4));
+  alloc_s2 (&q->m.s2, N);
+  alloc_s2 (&q->n.s2, N);
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1])
+#pragma acc parallel loop copy(p.m.s2->e[:N])
+      for (int j = 0; j < N; j++)
+	p.m.s2->e[j]++;
+#pragma acc exit data delete(p.m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->e[i] != 99)
+      abort ();
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1])
+#pragma acc enter data copyin(p.m.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.m.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    p.m.s2->f->a[j]++;
+	    p.m.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(p.m.s2->f[:1])
+#pragma acc exit data delete(p.m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->f->a[i] != 99 || p.m.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1]) copyin(p.n.s2[:1])
+#pragma acc enter data copyin(p.m.s2->f[:1]) copyin(p.n.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.m.s2->f->c[:N]) \
+			  copy(p.n.s2->f->a[:N]) copy(p.n.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    p.m.s2->f->a[j]++;
+	    p.m.s2->f->c[j]++;
+	    p.n.s2->f->a[j]++;
+	    p.n.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(p.m.s2->f[:1]) delete(p.n.s2->f[:1])
+#pragma acc exit data delete(p.m.s2[:1]) delete(p.n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->f->a[i] != 99 || p.m.s2->f->c[i] != 99
+	|| p.n.s2->f->a[i] != 99 || p.n.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (&p, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(p.m.s2[:1]) copyin(p.n.s2[:1])
+#pragma acc enter data copyin(p.n.s2->e[:N]) copyin(p.n.s2->f[:1]) \
+		       copyin(p.m.s2->f[:1])
+#pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.n.s2->f->a[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    p.m.s2->f->a[j]++;
+	    p.n.s2->f->a[j]++;
+	    p.n.s2->e[j]++;
+	  }
+#pragma acc exit data delete(p.m.s2->f[:1]) delete(p.n.s2->f[:1]) \
+		      copyout(p.n.s2->e[:N])
+#pragma acc exit data delete(p.m.s2[:1]) delete(p.n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (p.m.s2->f->a[i] != 99 || p.n.s2->f->a[i] != 99
+	|| p.n.s2->e[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1])
+#pragma acc parallel loop copy(q->m.s2->e[:N])
+      for (int j = 0; j < N; j++)
+	q->m.s2->e[j]++;
+#pragma acc exit data delete(q->m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->e[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1])
+#pragma acc enter data copyin(q->m.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->m.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    q->m.s2->f->a[j]++;
+	    q->m.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(q->m.s2->f[:1])
+#pragma acc exit data delete(q->m.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->f->a[i] != 99 || q->m.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1]) copyin(q->n.s2[:1])
+#pragma acc enter data copyin(q->m.s2->f[:1]) copyin(q->n.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->m.s2->f->c[:N]) \
+			  copy(q->n.s2->f->a[:N]) copy(q->n.s2->f->c[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    q->m.s2->f->a[j]++;
+	    q->m.s2->f->c[j]++;
+	    q->n.s2->f->a[j]++;
+	    q->n.s2->f->c[j]++;
+	  }
+#pragma acc exit data delete(q->m.s2->f[:1]) delete(q->n.s2->f[:1])
+#pragma acc exit data delete(q->m.s2[:1]) delete(q->n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->f->a[i] != 99 || q->m.s2->f->c[i] != 99
+	|| q->n.s2->f->a[i] != 99 || q->n.s2->f->c[i] != 99)
+      abort ();
+
+  zero_arrays (q, N);
+
+  for (int i = 0; i < 99; i++)
+    {
+#pragma acc enter data copyin(q->m.s2[:1]) copyin(q->n.s2[:1])
+#pragma acc enter data copyin(q->n.s2->e[:N]) copyin(q->m.s2->f[:1]) \
+		       copyin(q->n.s2->f[:1])
+#pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->n.s2->f->a[:N])
+	for (int j = 0; j < N; j++)
+	  {
+	    q->m.s2->f->a[j]++;
+	    q->n.s2->f->a[j]++;
+	    q->n.s2->e[j]++;
+	  }
+#pragma acc exit data delete(q->m.s2->f[:1]) delete(q->n.s2->f[:1]) \
+		      copyout(q->n.s2->e[:N])
+#pragma acc exit data delete(q->m.s2[:1]) delete(q->n.s2[:1])
+    }
+
+  for (i = 0; i < N; i++)
+    if (q->m.s2->f->a[i] != 99 || q->n.s2->f->a[i] != 99
+	|| q->n.s2->e[i] != 99)
+      abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c
new file mode 100644
index 0000000..a11c647
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c
@@ -0,0 +1,83 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <stdio.h>
+
+typedef struct {
+  int *a;
+  int *b;
+  int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  const int S = 32;
+  mystruct *m = (mystruct *) calloc (S, sizeof (*m));
+  int i, j;
+
+  for (i = 0; i < S; i++)
+    {
+      m[i].a = (int *) malloc (N * sizeof (int));
+      m[i].b = (int *) malloc (N * sizeof (int));
+      m[i].c = (int *) malloc (N * sizeof (int));
+    }
+
+  for (j = 0; j < S; j++)
+    for (i = 0; i < N; i++)
+      {
+	m[j].a[i] = 0;
+	m[j].b[i] = 0;
+	m[j].c[i] = 0;
+      }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j, k;
+      for (k = 0; k < S; k++)
+#pragma acc parallel loop copy(m[k].a[0:N])
+        for (j = 0; j < N; j++)
+          m[k].a[j]++;
+
+      for (k = 0; k < S; k++)
+#pragma acc parallel loop copy(m[k].b[0:N], m[k].c[5:N-10])
+	for (j = 0; j < N; j++)
+	  {
+	    m[k].b[j]++;
+	    if (j > 5 && j < N - 5)
+	      m[k].c[j]++;
+	}
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  for (j = 0; j < S; j++)
+    {
+      for (i = 0; i < N; i++)
+	{
+	  if (m[j].a[i] != 99)
+	    abort ();
+	  if (m[j].b[i] != 99)
+	    abort ();
+	  if (i > 5 && i < N-5)
+	    {
+	      if (m[j].c[i] != 99)
+		abort ();
+	    }
+	  else
+	    {
+	      if (m[j].c[i] != 0)
+		abort ();
+	    }
+	}
+
+      free (m[j].a);
+      free (m[j].b);
+      free (m[j].c);
+    }
+  free (m);
+
+  return 0;
+}