23baa717c9
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.
231 lines
5.5 KiB
C
231 lines
5.5 KiB
C
#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;
|
|
}
|