git
/
gcc
mirror of git://gcc.gnu.org/git/gcc.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

[og10] Rework indirect struct handling for OpenACC in gimplify.c 

This patch reworks indirect struct handling in gimplify.c (i.e. for
struct components mapped with "mystruct->a[0:n]", "mystruct->b", etc.),
for OpenACC.  The key observation leading to these changes was that
component mappings of references-to-structures is already implemented
and working, and indirect struct component handling via a pointer can
work quite similarly.  That lets us remove some earlier, special-case
handling for mapping indirect struct component accesses for OpenACC,
which required the pointed-to struct to be manually mapped before the
indirect component mapping.

With this patch, you can map struct components directly (e.g. an array
slice "mystruct->a[0:n]") just like you can map a non-indirect struct
component slice ("mystruct.a[0:n]"). Both references-to-pointers (with
the former syntax) and references to structs (with the latter syntax)
work now.

For Fortran class pointers, we no longer re-use GOMP_MAP_TO_PSET for the
class metadata (the structure that points to the class data and vptr)
-- it is instead treated as any other struct.

For C++, the struct handling also works for class members ("this->foo"),
without having to explicitly map "this[:1]" first.

For OpenACC, we permit chained indirect component references
("mystruct->a->b[0:n]"), though only the last part of such mappings will
trigger an attach/detach operation.  To properly use such a construct
on the target, you must still manually map "mystruct->a[:1]" first --
but there's no need to map "mystruct[:1]" explicitly before that.

This version of the patch avoids altering code paths for OpenMP,
where possible.

2021-05-19  Julian Brown  <julian@codesourcery.com>

gcc/fortran/
	* trans-openmp.c (gfc_trans_omp_clauses): Don't create GOMP_MAP_TO_PSET
	mappings for class metadata, nor GOMP_MAP_POINTER mappings for
	POINTER_TYPE_P decls.

gcc/
	* gimplify.c (extract_base_bit_offset): Add BASE_IND and OPENMP
	parameters.  Handle pointer-typed indirect references for OpenACC
	alongside reference-typed ones.
	(strip_components_and_deref, aggregate_base_p): New functions.
	(build_struct_group): Add pointer type indirect ref handling,
	including chained references, for OpenACC.  Also handle references to
	structs for OpenACC.  Conditionalise bits for OpenMP only where
	appropriate.
	(gimplify_scan_omp_clauses): Rework pointer-type indirect structure
	access handling to work more like the reference-typed handling for
	OpenACC only.
	* omp-low.c (scan_sharing_clauses): Handle pointer-type indirect struct
	references, and references to pointers to structs also.

gcc/testsuite/
	* g++.dg/goacc/member-array-acc.C: New test.
	* g++.dg/gomp/member-array-omp.C: New test.

libgomp/
	* 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.
This commit is contained in:
Julian Brown 2021-05-18 10:22:56 -07:00
parent 806fef8661
commit 091860deb6
8 changed files with 619 additions and 58 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

20
gcc/fortran/trans-openmp.c
View File

@ -2696,30 +2696,16 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
tree present = gfc_omp_check_optional_argument (decl, true);
if (openacc && n->sym->ts.type == BT_CLASS)
{
tree type = TREE_TYPE (decl);
if (n->sym->attr.optional)
sorry ("optional class parameter");
if (POINTER_TYPE_P (type))
{
node4 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER);
OMP_CLAUSE_DECL (node4) = decl;
OMP_CLAUSE_SIZE (node4) = size_int (0);
decl = build_fold_indirect_ref (decl);
}
tree ptr = gfc_class_data_get (decl);
ptr = build_fold_indirect_ref (ptr);
OMP_CLAUSE_DECL (node) = ptr;
OMP_CLAUSE_SIZE (node) = gfc_class_vtab_size_get (decl);
node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
OMP_CLAUSE_DECL (node2) = decl;
OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH_DETACH);
OMP_CLAUSE_DECL (node3) = gfc_class_data_get (decl);
OMP_CLAUSE_SIZE (node3) = size_int (0);
OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_ATTACH_DETACH);
OMP_CLAUSE_DECL (node2) = gfc_class_data_get (decl);
OMP_CLAUSE_SIZE (node2) = size_int (0);
goto finalize_map_clause;
}
else if (POINTER_TYPE_P (TREE_TYPE (decl))

215
gcc/gimplify.c
View File

@ -8323,8 +8323,9 @@ build_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end,
has array type, else return NULL. */
static tree
extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
poly_offset_int *poffsetp, tree *offsetp)
extract_base_bit_offset (tree base, tree *base_ind, tree *base_ref,
poly_int64 *bitposp, poly_offset_int *poffsetp,
tree *offsetp, bool openmp)
{
tree offset;
poly_int64 bitsize, bitpos;
@ -8332,6 +8333,9 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
int unsignedp, reversep, volatilep = 0;
poly_offset_int poffset;
if (base_ind)
*base_ind = NULL_TREE;
if (base_ref)
*base_ref = NULL_TREE;
@ -8352,14 +8356,29 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
base = get_inner_reference (base, &bitsize, &bitpos, &offset, &mode,
&unsignedp, &reversep, &volatilep);
tree orig_base = base;
if (!openmp
&& (TREE_CODE (base) == INDIRECT_REF
|| (TREE_CODE (base) == MEM_REF
&& integer_zerop (TREE_OPERAND (base, 1))))
&& TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == POINTER_TYPE)
{
if (base_ind)
*base_ind = base;
base = TREE_OPERAND (base, 0);
}
if ((TREE_CODE (base) == INDIRECT_REF
|| (TREE_CODE (base) == MEM_REF
&& integer_zerop (TREE_OPERAND (base, 1))))
&& DECL_P (TREE_OPERAND (base, 0))
&& TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == REFERENCE_TYPE)
base = TREE_OPERAND (base, 0);
{
if (base_ref)
*base_ref = base;
base = TREE_OPERAND (base, 0);
}
if (!openmp)
STRIP_NOPS (base);
if (offset && poly_int_tree_p (offset))
{
@ -8376,10 +8395,6 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
*poffsetp = poffset;
*offsetp = offset;
/* Set *BASE_REF if BASE was a dereferenced reference variable. */
if (base_ref && orig_base != base)
*base_ref = orig_base;
return base;
}
@ -8401,6 +8416,48 @@ is_or_contains_p (tree expr, tree base_ptr)
return operand_equal_p (expr, base_ptr);
}
/* Remove COMPONENT_REFS and indirections from EXPR. */
static tree
strip_components_and_deref (tree expr)
{
while (TREE_CODE (expr) == COMPONENT_REF
|| TREE_CODE (expr) == INDIRECT_REF
|| (TREE_CODE (expr) == MEM_REF
&& integer_zerop (TREE_OPERAND (expr, 1))))
expr = TREE_OPERAND (expr, 0);
return expr;
}
/* Return TRUE if EXPR is something we will use as the base of an aggregate
access, either:
- a DECL_P.
- a struct component with no indirection ("a.b.c").
- a struct component with indirection ("a->b->c").
*/
static bool
aggregate_base_p (tree expr)
{
while (TREE_CODE (expr) == COMPONENT_REF
&& (DECL_P (TREE_OPERAND (expr, 0))
|| (TREE_CODE (TREE_OPERAND (expr, 0)) == COMPONENT_REF)))
expr = TREE_OPERAND (expr, 0);
if (DECL_P (expr))
return true;
if (TREE_CODE (expr) == COMPONENT_REF
&& (TREE_CODE (TREE_OPERAND (expr, 0)) == INDIRECT_REF
|| (TREE_CODE (TREE_OPERAND (expr, 0)) == MEM_REF
&& integer_zerop (TREE_OPERAND (TREE_OPERAND (expr, 0), 1)))))
return true;
return false;
}
/* Implement OpenMP 5.x map ordering rules for target directives. There are
several rules, and with some level of ambiguity, hopefully we can at least
collect the complexity here in one place. */
@ -8788,10 +8845,12 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
{
poly_offset_int coffset;
poly_int64 cbitpos;
tree base_ref, tree_coffset;
tree base_ind, base_ref, tree_coffset;
bool openmp = !(region_type & ORT_ACC);
tree base = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ref,
&cbitpos, &coffset, &tree_coffset);
tree base = extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ind,
&base_ref, &cbitpos, &coffset,
&tree_coffset, openmp);
bool do_map_struct = (base == decl && !tree_coffset);
@ -8830,12 +8889,7 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
return NULL_TREE;
/* Nor for attach_detach for OpenMP. */
if ((code == OMP_TARGET
|| code == OMP_TARGET_DATA
|| code == OMP_TARGET_UPDATE
|| code == OMP_TARGET_ENTER_DATA
|| code == OMP_TARGET_EXIT_DATA)
&& attach_detach)
if (openmp && attach_detach)
{
if (DECL_P (decl))
{
@ -8858,12 +8912,15 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
OMP_CLAUSE_SET_MAP_KIND (l, k);
if (base_ref)
if (!openmp && base_ind)
OMP_CLAUSE_DECL (l) = unshare_expr (base_ind);
else if (base_ref)
OMP_CLAUSE_DECL (l) = unshare_expr (base_ref);
else
{
OMP_CLAUSE_DECL (l) = unshare_expr (decl);
if (!DECL_P (OMP_CLAUSE_DECL (l))
if (openmp
&& !DECL_P (OMP_CLAUSE_DECL (l))
&& (gimplify_expr (&OMP_CLAUSE_DECL (l), pre_p, NULL,
is_gimple_lvalue, fb_lvalue) == GS_ERROR))
return error_mark_node;
@ -8919,6 +8976,48 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l);
OMP_CLAUSE_CHAIN (l) = c2;
}
else if (!openmp
&& (base_ind || base_ref)
&& (region_type & ORT_TARGET))
{
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
enum gomp_map_kind mkind = base_ref ? GOMP_MAP_FIRSTPRIVATE_REFERENCE
: GOMP_MAP_FIRSTPRIVATE_POINTER;
OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
OMP_CLAUSE_SIZE (c2) = size_zero_node;
tree sdecl = strip_components_and_deref (decl);
if (DECL_P (decl)
&& (POINTER_TYPE_P (TREE_TYPE (sdecl))
|| TREE_CODE (TREE_TYPE (sdecl)) == REFERENCE_TYPE))
{
/* Insert after struct node. */
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (l);
OMP_CLAUSE_DECL (c2) = decl;
OMP_CLAUSE_CHAIN (l) = c2;
}
else
{
/* If the ultimate base for this component access is not a
pointer or reference, that means it is a struct component
access itself. Insert a node to be processed on the next
iteration of our caller's loop, which will subsequently be
turned into a new GOMP_MAP_STRUCT mapping itself.
We need to do this else the non-DECL_P base won't be
rewritten correctly in the offloaded region. */
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FORCE_PRESENT);
OMP_CLAUSE_DECL (c2) = unshare_expr (decl);
OMP_CLAUSE_SIZE (c2) = (DECL_P (decl)
? DECL_SIZE_UNIT (decl)
: TYPE_SIZE_UNIT (TREE_TYPE (decl)));
tree *next_node = &OMP_CLAUSE_CHAIN (*list_p);
OMP_CLAUSE_CHAIN (c2) = *next_node;
*next_node = c2;
return NULL_TREE;
}
}
*flags = GOVD_MAP | GOVD_EXPLICIT;
if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr || attach_detach)
*flags |= GOVD_SEEN;
@ -8927,7 +9026,8 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
/* If this is a *pointer-to-struct expression, make sure a
firstprivate map of the base-pointer exists. */
if (component_ref_p
if (openmp
&& component_ref_p
&& ((TREE_CODE (decl) == MEM_REF
&& integer_zerop (TREE_OPERAND (decl, 1)))
|| INDIRECT_REF_P (decl))
@ -8957,10 +9057,12 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
n->value |= GOVD_SEEN;
sc = &OMP_CLAUSE_CHAIN (*osc);
/* The struct mapping might be immediately followed by a
FIRSTPRIVATE_REFERENCE if it is a reference. (This added node is
removed in omp-low.c after it has been processed there.) */
FIRSTPRIVATE_POINTER and/or FIRSTPRIVATE_REFERENCE -- if it's an
indirect access or a reference, or both. (This added node is removed
in omp-low.c after it has been processed there.) */
if (*sc != c
&& OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
&& (OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
sc = &OMP_CLAUSE_CHAIN (*sc);
for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc))
if ((ptr || attach_detach) && sc == prev_list_p)
@ -8975,9 +9077,10 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
poly_offset_int offset;
poly_int64 bitpos;
tree tree_offset;
tree base = extract_base_bit_offset (sc_decl, NULL, &bitpos,
&offset, &tree_offset);
if (base != decl)
tree base = extract_base_bit_offset (sc_decl, NULL, NULL,
&bitpos, &offset,
&tree_offset, openmp);
if (!base || !operand_equal_p (base, decl, 0))
break;
if (scp)
continue;
@ -9098,8 +9201,9 @@ build_struct_group (struct gimplify_omp_ctx *ctx,
}
else if (*sc != c)
{
if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
== GS_ERROR)
if (openmp
&& (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
== GS_ERROR))
return error_mark_node;
/* In the non-pointer case, the mapping clause itself is moved into
the correct position in the struct component list, which in this
@ -9702,10 +9806,43 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
tree indir_base = NULL_TREE;
tree orig_decl = decl;
tree decl_ref = NULL_TREE;
if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
&& TREE_CODE (*pd) == COMPONENT_REF
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
&& code != OACC_UPDATE)
if ((region_type & ORT_ACC) && TREE_CODE (decl) == COMPONENT_REF)
{
/* Strip off component refs from RHS of e.g. "a->b->c.d.e"
(which would leave "a->b" in that case). This is intended
to be equivalent to the base finding done by
get_inner_reference. */
while (TREE_CODE (decl) == COMPONENT_REF
&& (DECL_P (TREE_OPERAND (decl, 0))
|| (TREE_CODE (TREE_OPERAND (decl, 0))
== COMPONENT_REF)))
decl = TREE_OPERAND (decl, 0);
if (TREE_CODE (decl) == COMPONENT_REF)
decl = TREE_OPERAND (decl, 0);
/* Strip off RHS from "a->b". */
if ((TREE_CODE (decl) == INDIRECT_REF
|| (TREE_CODE (decl) == MEM_REF
&& integer_zerop (TREE_OPERAND (decl, 1))))
&& (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
== POINTER_TYPE))
decl = TREE_OPERAND (decl, 0);
/* Strip off RHS from "a_ref.b" (where a_ref is
reference-typed). */
if (TREE_CODE (decl) == INDIRECT_REF
&& DECL_P (TREE_OPERAND (decl, 0))
&& (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
== REFERENCE_TYPE))
decl = TREE_OPERAND (decl, 0);
STRIP_NOPS (decl);
}
else if ((region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
&& TREE_CODE (*pd) == COMPONENT_REF
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
&& code != OACC_UPDATE)
{
while (TREE_CODE (decl) == COMPONENT_REF)
{
@ -9808,11 +9945,13 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
if (code == OACC_UPDATE
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
if ((DECL_P (decl)
|| (component_ref_p
&& (INDIRECT_REF_P (decl)
|| TREE_CODE (decl) == MEM_REF
|| TREE_CODE (decl) == ARRAY_REF)))
if ((((region_type & ORT_ACC) && aggregate_base_p (decl))
|| (!(region_type & ORT_ACC)
&& (DECL_P (decl)
|| (component_ref_p
&& (INDIRECT_REF_P (decl)
|| TREE_CODE (decl) == MEM_REF
|| TREE_CODE (decl) == ARRAY_REF)))))
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH

16
gcc/omp-low.c
View File

@ -1651,8 +1651,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
if (TREE_CODE (decl) == COMPONENT_REF
|| (TREE_CODE (decl) == INDIRECT_REF
&& TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
&& (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
== REFERENCE_TYPE)))
&& (((TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
== REFERENCE_TYPE)
|| (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
== POINTER_TYPE)))))
break;
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
@ -13504,6 +13506,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
is_ref = false;
bool ref_to_array = false;
bool ref_to_ptr = false;
if (is_ref)
{
type = TREE_TYPE (type);
@ -13522,6 +13525,12 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
new_var = decl2;
type = TREE_TYPE (new_var);
}
else if (TREE_CODE (type) == REFERENCE_TYPE
&& TREE_CODE (TREE_TYPE (type)) == POINTER_TYPE)
{
type = TREE_TYPE (type);
ref_to_ptr = true;
}
x = build_receiver_ref (prev, false, ctx);
x = fold_convert_loc (clause_loc, type, x);
if (!integer_zerop (OMP_CLAUSE_SIZE (c)))
@ -13544,7 +13553,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (ref_to_array)
x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
if (is_ref && !ref_to_array)
if ((is_ref && !ref_to_array)
|| ref_to_ptr)
{
tree t = create_tmp_var_raw (type, get_name (var));
gimple_add_tmp_var (t);

13
gcc/testsuite/g++.dg/goacc/member-array-acc.C Normal file
View File

@ -0,0 +1,13 @@
/* { dg-do compile } */
/* { dg-additional-options "-fdump-tree-gimple" } */
struct Foo {
float *a;
void init(int N) {
a = new float[N];
#pragma acc enter data create(a[0:N])
}
};
int main() { Foo x; x.init(1024); }
/* { dg-final { scan-tree-dump {struct:\*\(struct Foo \*\) this \[len: 1\]\) map\(alloc:\(\(struct Foo \*\) this\)->a \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:\(\(struct Foo \*\) this\)->a \[bias: 0\]\)} "gimple" } } */

13
gcc/testsuite/g++.dg/gomp/member-array-omp.C Normal file
View File

@ -0,0 +1,13 @@
/* { dg-do compile } */
/* { dg-additional-options "-fdump-tree-gimple" } */
struct Foo {
float *a;
void init(int N) {
a = new float[N];
#pragma omp target enter data map(alloc:a[0:N])
}
};
int main() { Foo x; x.init(1024); }
/* { dg-final { scan-tree-dump {map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:this->a \[bias: 0\]\)} "gimple" } } */

101
libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C Normal file
View File

@ -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;
}

68
libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c Normal file
View File

@ -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;
}

231
libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c Normal file
View File

@ -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;
}