mirror of
git://gcc.gnu.org/git/gcc.git
synced 2025-04-05 05:20:24 +08:00
OpenMP/OpenACC struct sibling list gimplification extension and rework
This patch refactors struct sibling-list processing in gimplify.cc, and adjusts some related mapping-clause processing in the Fortran FE and omp-low.cc accordingly. 2022-09-13 Julian Brown <julian@codesourcery.com> gcc/fortran/ * trans-openmp.cc (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.cc (gimplify_omp_var_data): Remove GOVD_MAP_HAS_ATTACHMENTS. (GOMP_FIRSTPRIVATE_IMPLICIT): Renumber. (insert_struct_comp_map): Refactor function into... (build_omp_struct_comp_nodes): This new function. Remove list handling and improve self-documentation. (extract_base_bit_offset): Remove BASE_REF, OFFSETP parameters. Move code to strip outer parts of address out of function, but strip no-op conversions. (omp_mapping_group): Add DELETED field for use during reindexing. (omp_strip_components_and_deref, omp_strip_indirections): New functions. (omp_group_last, omp_group_base): Add GOMP_MAP_STRUCT handling. (omp_gather_mapping_groups): Initialise DELETED field for new groups. (omp_index_mapping_groups): Notice DELETED groups when (re)indexing. (omp_siblist_insert_node_after, omp_siblist_move_node_after, omp_siblist_move_nodes_after, omp_siblist_move_concat_nodes_after): New helper functions. (omp_accumulate_sibling_list): New function to build up GOMP_MAP_STRUCT node groups for sibling lists. Outlined from gimplify_scan_omp_clauses. (omp_build_struct_sibling_lists): New function. (gimplify_scan_omp_clauses): Remove struct_map_to_clause, struct_seen_clause, struct_deref_set. Call omp_build_struct_sibling_lists as pre-pass instead of handling sibling lists in the function's main processing loop. (gimplify_adjust_omp_clauses_1): Remove GOVD_MAP_HAS_ATTACHMENTS handling, unused now. * omp-low.cc (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. * g++.dg/gomp/target-3.C: Update expected output. * g++.dg/gomp/target-lambda-1.C: Likewise. * g++.dg/gomp/target-this-2.C: Likewise. * c-c++-common/goacc/deep-copy-arrayofstruct.c: Move test from here. * c-c++-common/gomp/target-50.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. * testsuite/libgomp.oacc-c-c++-common/deep-copy-arrayofstruct.c: Move test to here, make "run" test.
This commit is contained in:
Julian Brown
parent
cd14c97cd9
commit
23baa717c9
@ -3125,30 +3125,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))
|
||||
|
||||
1514
gcc/gimplify.cc
1514
gcc/gimplify.cc
File diff suppressed because it is too large
Load Diff
@ -1636,8 +1636,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)
|
||||
@ -14015,6 +14017,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);
|
||||
@ -14033,6 +14036,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 (OMP_CLAUSE_DECL (prev), false, ctx);
|
||||
x = fold_convert_loc (clause_loc, type, x);
|
||||
if (!integer_zerop (OMP_CLAUSE_SIZE (c)))
|
||||
@ -14049,7 +14058,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);
|
||||
|
||||
23
gcc/testsuite/c-c++-common/gomp/target-50.c
Normal file
23
gcc/testsuite/c-c++-common/gomp/target-50.c
Normal file
@ -0,0 +1,23 @@
|
||||
/* { dg-do compile } */
|
||||
/* { dg-additional-options "-fdump-tree-gimple" } */
|
||||
|
||||
typedef struct
|
||||
{
|
||||
int *arr;
|
||||
} L;
|
||||
|
||||
int main()
|
||||
{
|
||||
L *tmp;
|
||||
|
||||
/* There shouldn't be an order dependency here... */
|
||||
|
||||
#pragma omp target map(to: tmp->arr) map(tofrom: tmp->arr[0:10])
|
||||
{ }
|
||||
|
||||
#pragma omp target map(tofrom: tmp->arr[0:10]) map(to: tmp->arr)
|
||||
{ }
|
||||
/* { dg-final { scan-tree-dump-times {map\(struct:\*tmp \[len: 1\]\) map\(alloc:tmp[._0-9]*->arr \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(attach:tmp[._0-9]*->arr \[bias: 0\]\)} 2 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */
|
||||
|
||||
return 0;
|
||||
}
|
||||
13
gcc/testsuite/g++.dg/goacc/member-array-acc.C
Normal file
13
gcc/testsuite/g++.dg/goacc/member-array-acc.C
Normal 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:this->a \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:this->a \[bias: 0\]\)} "gimple" } } */
|
||||
13
gcc/testsuite/g++.dg/gomp/member-array-omp.C
Normal file
13
gcc/testsuite/g++.dg/gomp/member-array-omp.C
Normal 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" } } */
|
||||
@ -33,4 +33,6 @@ T<N>::bar (int x)
|
||||
|
||||
template struct T<0>;
|
||||
|
||||
/* { dg-final { scan-tree-dump-times "map\\(struct:\\*this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "map\\(struct:\\*\\(struct S \\*\\) this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
|
||||
|
||||
/* { dg-final { scan-tree-dump-times "map\\(struct:\\*\\(struct T \\*\\) this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
|
||||
|
||||
@ -87,8 +87,7 @@ int main (void)
|
||||
return 0;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)
|
||||
} "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */
|
||||
|
||||
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */
|
||||
|
||||
|
||||
@ -46,4 +46,4 @@ int main (void)
|
||||
return 0;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) firstprivate\(m\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:v \[len: [0-9]+\]\)} "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump {map\(alloc:MEM\[\(char \*\)_[0-9]+\] \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) firstprivate\(m\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:v \[len: [0-9]+\]\)} "gimple" } } */
|
||||
|
||||
101
libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C
Normal file
101
libgomp/testsuite/libgomp.oacc-c++/deep-copy-17.C
Normal 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
68
libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c
Normal 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
231
libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c
Normal 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;
|
||||
}
|
||||
@ -1,4 +1,4 @@
|
||||
/* { dg-do compile } */
|
||||
/* { dg-do run } */
|
||||
|
||||
#include <stdlib.h>
|
||||
#include <stdio.h>
|
||||
Loading…
x
Reference in New Issue
Block a user