commit 56101feb78bc2e3344159f96b7d0ab9eaf4bd529
Author: Julian Brown <julian@codesourcery.com>
Date:   Wed Jan 30 04:54:24 2019 -0800

    [og8] Attach/detach array slices on dereferenced struct members
    
            gcc/c/
            * c-typeck.c (handle_omp_array_sections_1): Handle chained dereferences.
            (c_finish_omp_clauses): Likewise.
    
            gcc/cp/
            * semantics.c (handle_omp_array_sections_1): Handle array section on
            dereferenced struct member.
            (finish_omp_clauses): Don't error on multiple dereferenced struct
            elements with the same base.
    
            gcc/
            * gimplify.c (gimplify_scan_omp_clauses): Handle array sections on
            dereferenced struct members.
    
            gcc/testsuite/
            * c-c++-common/goacc/deep-copy-arrayofstruct.c: New test.
    
            libgomp/
            * testsuite/libgomp.oacc-c++/deep-copy-12.C: New test.
            * testsuite/libgomp.oacc-c++/deep-copy-13.C: New test.
            * testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c: New test.
            * testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c: New test.
            * testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c: New test.
            * testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c: New test.

diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index d25e2d8c14c..7f021649216 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12445,9 +12445,16 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		  return error_mark_node;
 		}
 	      t = TREE_OPERAND (t, 0);
+	      if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
+		{
+		  if (maybe_ne (mem_ref_offset (t), 0))
+	            error_at (OMP_CLAUSE_LOCATION (c),
+			      "cannot dereference %qE in %qs clause", t,
+			      omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		  else
+		    t = TREE_OPERAND (t, 0);
+		}
 	    }
-	  if (TREE_CODE (t) == MEM_REF)
-	    t = TREE_OPERAND (t, 0);
 	}
       if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
 	{
@@ -13750,11 +13757,18 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		      break;
 		    }
 		  t = TREE_OPERAND (t, 0);
+		  if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
+	            {
+		      if (maybe_ne (mem_ref_offset (t), 0))
+			error_at (OMP_CLAUSE_LOCATION (c),
+				  "cannot dereference %qE in %qs clause", t,
+				  omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		      else
+			t = TREE_OPERAND (t, 0);
+		    }
 		}
 	      if (remove)
 		break;
-	      if (TREE_CODE (t) == MEM_REF)
-		t = TREE_OPERAND (t, 0);
 	      if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
 		{
 		  if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 497fd39b10c..72c4dcec2b3 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -4557,6 +4557,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		  return error_mark_node;
 		}
 	      t = TREE_OPERAND (t, 0);
+	      if (ort == C_ORT_ACC && TREE_CODE (t) == INDIRECT_REF)
+	        t = TREE_OPERAND (t, 0);
 	    }
 	  if (REFERENCE_REF_P (t))
 	    t = TREE_OPERAND (t, 0);
@@ -6941,7 +6943,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      else
 		bitmap_set_bit (&generic_head, DECL_UID (t));
 	    }
-	  else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+	  else if (bitmap_bit_p (&map_head, DECL_UID (t))
+		   && (ort != C_ORT_ACC
+		       || !bitmap_bit_p (&map_field_head, DECL_UID (t))))
 	    {
 	      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 		error ("%qD appears more than once in motion clauses", t);
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index a6a4d2a68dd..8bf11eb659e 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -7808,6 +7808,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
   struct gimplify_omp_ctx *ctx, *outer_ctx;
   tree c;
   hash_map<tree, tree> *struct_map_to_clause = NULL;
+  hash_set<tree> *struct_deref_set = NULL;
   tree *prev_list_p = NULL;
 
   ctx = new_omp_context (region_type);
@@ -8211,7 +8212,35 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  pd = &TREE_OPERAND (decl, 0);
 		  decl = TREE_OPERAND (decl, 0);
 		}
-	      if (TREE_CODE (decl) == COMPONENT_REF)
+	      bool indir_p = false;
+	      tree orig_decl = decl;
+	      tree decl_ref = NULL_TREE;
+	      if ((region_type & ORT_ACC) != 0
+		  && TREE_CODE (*pd) == COMPONENT_REF
+		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER
+		  && code != OACC_UPDATE)
+		{
+		  while (TREE_CODE (decl) == COMPONENT_REF)
+		    {
+		      decl = TREE_OPERAND (decl, 0);
+		      if ((TREE_CODE (decl) == MEM_REF
+			   && integer_zerop (TREE_OPERAND (decl, 1)))
+			  || INDIRECT_REF_P (decl))
+			{
+			  indir_p = true;
+			  decl = TREE_OPERAND (decl, 0);
+			}
+		      if (TREE_CODE (decl) == INDIRECT_REF
+			  && DECL_P (TREE_OPERAND (decl, 0))
+			  && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+			      == REFERENCE_TYPE))
+			{
+			  decl_ref = decl;
+			  decl = TREE_OPERAND (decl, 0);
+			}
+		    }
+		}
+	      else if (TREE_CODE (decl) == COMPONENT_REF)
 		{
 		  while (TREE_CODE (decl) == COMPONENT_REF)
 		    decl = TREE_OPERAND (decl, 0);
@@ -8221,6 +8250,52 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			  == REFERENCE_TYPE))
 		    decl = TREE_OPERAND (decl, 0);
 		}
+	      if (decl != orig_decl && DECL_P (decl) && indir_p)
+		{
+		  gomp_map_kind k = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
+							     : GOMP_MAP_ATTACH;
+		  /* We have a dereference of a struct member.  Make this an
+		     attach/detach operation, and ensure the base pointer is
+		     mapped as a FIRSTPRIVATE_POINTER.  */
+		  OMP_CLAUSE_SET_MAP_KIND (c, k);
+		  flags = GOVD_MAP | GOVD_SEEN | GOVD_EXPLICIT;
+		  if (k == GOMP_MAP_ATTACH
+		      && (!struct_deref_set
+			  || !struct_deref_set->contains (decl)))
+		    {
+		      if (!struct_deref_set)
+		        struct_deref_set = new hash_set<tree> ();
+		      /* As well as the attach, we also need a
+			 FIRSTPRIVATE_POINTER clause to properly map the
+			 pointer to the struct base.  */
+		      tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+						  OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC);
+		      OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c2)
+			= 1;
+		      tree charptr_zero
+		        = build_int_cst (build_pointer_type (char_type_node),
+					 0);
+		      OMP_CLAUSE_DECL (c2)
+			= build2 (MEM_REF, char_type_node,
+				  decl_ref ? decl_ref : decl, charptr_zero);
+		      OMP_CLAUSE_SIZE (c2) = size_zero_node;
+		      tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+						  OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_SET_MAP_KIND (c3,
+					       GOMP_MAP_FIRSTPRIVATE_POINTER);
+		      OMP_CLAUSE_DECL (c3) = decl;
+		      OMP_CLAUSE_SIZE (c3) = size_zero_node;
+		      *list_p = c2;
+		      OMP_CLAUSE_CHAIN (c2) = c3;
+		      OMP_CLAUSE_CHAIN (c3) = c;
+		      c = c3;
+		      list_p = &OMP_CLAUSE_CHAIN (c3);
+
+		      struct_deref_set->add (decl);
+		    }
+		  goto do_add_decl;
+		}
 	      if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
 		  == GS_ERROR)
 		{
@@ -8831,6 +8906,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
   gimplify_omp_ctxp = ctx;
   if (struct_map_to_clause)
     delete struct_map_to_clause;
+  if (struct_deref_set)
+    delete struct_deref_set;
 }
 
 /* Return true if DECL is a candidate for shared to firstprivate
diff --git a/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c
new file mode 100644
index 00000000000..d411bcfa8e7
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c
@@ -0,0 +1,84 @@
+/* { dg-do compile } */
+
+#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]) /* { dg-error "expected .\\\). before .\\\.. token" } */
+        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]) /* { dg-error "expected .\\\). before .\\\.. token" } */
+	/* { dg-error ".m. appears more than once in data clauses" "" { target c++ } .-1 } */
+	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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C
new file mode 100644
index 00000000000..771876afd2d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C
@@ -0,0 +1,72 @@
+#include <stdlib.h>
+
+/* Test attach/detach with dereferences of reference to pointer to struct.  */
+
+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));
+  mystruct *&mref = m;
+  int i;
+
+  mref->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++)
+    {
+      mref->a[i] = 0;
+      m->b[i] = 0;
+      m->c[i] = 0;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+#pragma acc parallel loop copy(mref->a[0:N])
+      for (j = 0; j < N; j++)
+        mref->a[j]++;
+#pragma acc parallel loop copy(mref->b[0:N], m->c[5:N-10])
+      for (j = 0; j < N; j++)
+	{
+	  mref->b[j]++;
+	  if (j > 5 && j < N - 5)
+	    m->c[j]++;
+	}
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  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++/deep-copy-13.C b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C
new file mode 100644
index 00000000000..98cf450c61d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C
@@ -0,0 +1,72 @@
+#include <stdlib.h>
+
+/* Test array slice with reference to pointer.  */
+
+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;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+      int *&ptr = m->a;
+#pragma acc parallel loop copy(ptr[0:N])
+      for (j = 0; j < N; j++)
+        ptr[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]++;
+	}
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  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-10.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c
new file mode 100644
index 00000000000..37cde4ef059
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c
@@ -0,0 +1,53 @@
+#include <stdlib.h>
+
+/* Test asyncronous attach and detach operation.  */
+
+typedef struct {
+  int *a;
+  int *b;
+} mystruct;
+
+int
+main (int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct m;
+  int i;
+
+  m.a = (int *) malloc (N * sizeof (int));
+  m.b = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      m.a[i] = 0;
+      m.b[i] = 0;
+    }
+
+#pragma acc enter data copyin(m)
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+#pragma acc parallel loop copy(m.a[0:N]) async(i % 2)
+      for (j = 0; j < N; j++)
+        m.a[j]++;
+#pragma acc parallel loop copy(m.b[0:N]) async((i + 1) % 2)
+      for (j = 0; j < N; j++)
+        m.b[j]++;
+    }
+
+#pragma acc exit data copyout(m) wait(0, 1)
+
+  for (i = 0; i < N; i++)
+    {
+      if (m.a[i] != 99)
+        abort ();
+      if (m.b[i] != 99)
+        abort ();
+    }
+
+  free (m.a);
+  free (m.b);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c
new file mode 100644
index 00000000000..ed8ddcc54fa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c
@@ -0,0 +1,72 @@
+#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;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+
+  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]++;
+	}
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  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-14.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c
new file mode 100644
index 00000000000..04c70aed7b5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c
@@ -0,0 +1,63 @@
+#include <openacc.h>
+#include <stdlib.h>
+
+/* Test attach/detach operation with chained dereferences.  */
+
+typedef struct mystruct {
+  int *a;
+  struct mystruct *next;
+} 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->next = (mystruct *) malloc (sizeof (*m));
+  m->next->a = (int *) malloc (N * sizeof (int));
+  m->next->next = NULL;
+
+  for (i = 0; i < N; i++)
+    {
+      m->a[i] = 0;
+      m->next->a[i] = 0;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+  acc_copyin (m->next, sizeof (*m));
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+      acc_copyin (m->next->a, N * sizeof (int));
+      acc_attach ((void **) &m->next);
+      /* This will attach only the innermost pointer, i.e. "a[0:N]".  That's
+	 why we have to attach the "m->next" pointer manually above.  */
+#pragma acc parallel loop copy(m->next->a[0:N])
+      for (j = 0; j < N; j++)
+        m->next->a[j]++;
+      acc_detach ((void **) &m->next);
+      acc_copyout (m->next->a, N * sizeof (int));
+    }
+
+  acc_copyout (m->next, sizeof (*m));
+#pragma acc exit data copyout(m[0:1])
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 0)
+        abort ();
+      if (m->next->a[i] != 99)
+        abort ();
+    }
+
+  free (m->next->a);
+  free (m->next);
+  free (m->a);
+  free (m);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c
new file mode 100644
index 00000000000..28535c9e281
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c
@@ -0,0 +1,53 @@
+#include <stdlib.h>
+
+typedef struct {
+  int *a;
+  int *b;
+} 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));
+
+  for (i = 0; i < N; i++)
+    {
+      m->a[i] = 0;
+      m->b[i] = 0;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+      int *ptr = m->a;
+#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]) 
+      for (j = 0; j < N; j++)
+	m->b[j]++;
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 99)
+        abort ();
+      if (m->b[i] != 99)
+        abort ();
+    }
+
+  free (m->a);
+  free (m->b);
+  free (m);
+
+  return 0;
+}
