Hi,
This patch adds support for array slices on dereferenced struct
members, e.g.:
#pragma acc parallel copy(mystruct->a[0:n])
This works by making a new mapping pair for each struct pointer used in
the directive ("alloc(mystruct) firstprivate_pointer(mystruct)").
The C/C++ parsers permit chained dereferences
("mystruct->anotherstruct->bla[0:n]"). In this case, the current
implementation performs an attach/detach operation on the
final/innermost dereference only (so, "bla[0:n]" attaches to
the appropriate offset in "anotherstruct"). Other options might be to
explicitly disallow chained dereferences, or attach the whole chain
sequentially. The standard isn't helpful here (as of 2.6), but I think
that the chosen behaviour is reasonably consistent.
Arrays of structures aren't (yet?) supported
(either "copy(structarr[i].a[0:n])" or "copy(structarr[i]->a[0:n])"). I
added a basic test case for that.
Tested with offloading to nvptx, no regressions and the new tests pass.
I will apply shortly (to the og8 branch).
Thanks,
Julian
ChangeLog
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.
commit 56101feb78bc2e3344159f96b7d0ab9eaf4bd529
Author: Julian Brown <jul...@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;
+}
