From: Julian Brown Date: Thu, 31 Jan 2019 16:39:15 +0000 (-0800) Subject: [og8] Attach/detach array slices on dereferenced struct members X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=a7e6117e36b4ddf7cc6611196f39ee0960d32856;p=thirdparty%2Fgcc.git [og8] Attach/detach array slices on dereferenced struct members A couple of fixes folded in relative to og8 version. 2019-01-31 Julian Brown 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. (cherry picked from openacc-gcc-9-branch commit ad36f2e3a8ad481c6bbd6376a1ab25fce1da4d8d) --- diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index b923fc36c829..f8f5587779ae 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,8 @@ +2019-01-31 Julian Brown + + * gimplify.c (gimplify_scan_omp_clauses): Handle array sections on + dereferenced struct members. + 2018-12-14 Julian Brown * gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_HAS_ATTACHMENTS. diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp index e05fe9abba2c..dfedd90e1dc8 100644 --- a/gcc/c/ChangeLog.omp +++ b/gcc/c/ChangeLog.omp @@ -1,3 +1,8 @@ +2019-01-31 Julian Brown + + * c-typeck.c (handle_omp_array_sections_1): Handle chained dereferences. + (c_finish_omp_clauses): Likewise. + 2018-12-14 Julian Brown * c-parser.c (c_parser_omp_clause_name): Add parsing of attach and diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index d3383db66283..f741405f8560 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -12869,15 +12869,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec &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, - c_omp_map_clause_name (c, true)); - else - 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 (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) @@ -14454,18 +14454,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 (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, - c_omp_map_clause_name (c, true)); - else - 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/ChangeLog.omp b/gcc/cp/ChangeLog.omp index 5351c493d084..e397fab0df2c 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,3 +1,10 @@ +2019-01-31 Julian Brown + + * 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. + 2018-12-14 Julian Brown * parser.c (cp_parser_omp_clause_name): Support attach and detach diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index b7701701ead0..7fbff82abd64 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -4675,6 +4675,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec &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); @@ -7324,7 +7326,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_at (OMP_CLAUSE_LOCATION (c), diff --git a/gcc/gimplify.c b/gcc/gimplify.c index eb67d1a79cb5..79d0535c50a8 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8210,6 +8210,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, struct gimplify_omp_ctx *ctx, *outer_ctx; tree c; hash_map *struct_map_to_clause = NULL; + hash_set *struct_deref_set = NULL; tree *prev_list_p = NULL; int handled_depend_iterators = -1; int nowait = -1; @@ -8654,7 +8655,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); @@ -8664,6 +8693,51 @@ 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 (); + /* 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; + tree mapgrp = *prev_list_p; + *prev_list_p = c2; + OMP_CLAUSE_CHAIN (c3) = mapgrp; + OMP_CLAUSE_CHAIN (c2) = c3; + + struct_deref_set->add (decl); + } + goto do_add_decl; + } if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue) == GS_ERROR) { @@ -9385,6 +9459,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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 2c618552552d..33af5c50689f 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,7 @@ +2019-01-31 Julian Brown + + * c-c++-common/goacc/deep-copy-arrayofstruct.c: New test. + 2018-12-14 Julian Brown * c-c++-common/goacc/mdc-1.c: New test. 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 000000000000..d411bcfa8e7c --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c @@ -0,0 +1,84 @@ +/* { dg-do compile } */ + +#include +#include + +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/gcc/testsuite/gfortran.dg/goacc/derived-types-2.f90 b/gcc/testsuite/gfortran.dg/goacc/derived-types-2.f90 new file mode 100644 index 000000000000..d01583fac890 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/derived-types-2.f90 @@ -0,0 +1,14 @@ +module bar + type :: type1 + real(8), pointer, public :: p(:) => null() + end type + type :: type2 + class(type1), pointer :: p => null() + end type +end module + +subroutine foo (var) + use bar + type(type2), intent(inout) :: var + !$acc enter data create(var%p%p) +end subroutine diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 3c0bdf9a0992..7ee891413711 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,12 @@ +2019-01-31 Julian Brown + + * 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. + 2018-11-28 Julian Brown * libgomp.h (RC_CHECKING): New macro, disabled by default, guarding all 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 000000000000..771876afd2d2 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C @@ -0,0 +1,72 @@ +#include + +/* 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 000000000000..98cf450c61da --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C @@ -0,0 +1,72 @@ +#include + +/* 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 000000000000..37cde4ef059e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c @@ -0,0 +1,53 @@ +#include + +/* 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 000000000000..ed8ddcc54fa4 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c @@ -0,0 +1,72 @@ +#include + +/* 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 000000000000..04c70aed7b5a --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c @@ -0,0 +1,63 @@ +#include +#include + +/* 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 000000000000..28535c9e2819 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c @@ -0,0 +1,53 @@ +#include + +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; +}