This patch introduces a couple of compiler tests for the OpenACC attach and detach clauses.
I've committed it to openacc-gcc-8-branch. Cesar
2018-10-30 Cesar Philippidis <ce...@codesourcery.com> gcc/testsuite/ * c-c++-common/goacc/mdc-1.c: New test. * c-c++-common/goacc/mdc-2.c: New test. * g++.dg/goacc/mdc.C: New test. --- gcc/testsuite/c-c++-common/goacc/mdc-1.c | 54 +++++++++++++++++++ gcc/testsuite/c-c++-common/goacc/mdc-2.c | 62 +++++++++++++++++++++ gcc/testsuite/g++.dg/goacc/mdc.C | 68 ++++++++++++++++++++++++ 3 files changed, 184 insertions(+) create mode 100644 gcc/testsuite/c-c++-common/goacc/mdc-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/mdc-2.c create mode 100644 gcc/testsuite/g++.dg/goacc/mdc.C diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c new file mode 100644 index 00000000000..c20b94ddbdc --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c @@ -0,0 +1,54 @@ +/* Test OpenACC's support for manual deep copy, including the attach + and detach clauses. */ + +/* { dg-additional-options "-fdump-tree-omplower" } */ + +void +t1 () +{ + struct foo { + int *a, *b, c, d, *e; + } s; + + int *a, *z; + +#pragma acc enter data copyin(s) + { +#pragma acc data copy(s.a[0:10]) copy(z[0:10]) + { + s.e = z; +#pragma acc parallel loop attach(s.e) + for (int i = 0; i < 10; i++) + s.a[i] = s.e[i]; + + + a = s.e; +#pragma acc enter data attach(a) +#pragma acc exit data detach(a) + } + +#pragma acc enter data copyin(a) +#pragma acc acc enter data attach(s.e) +#pragma acc exit data detach(s.e) + +#pragma acc data attach(s.e) + { + } +#pragma acc exit data delete(a) + +#pragma acc exit data detach(a) finalize +#pragma acc exit data detach(s.a) finalize + } +} + +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.always_pointer:s.a .pointer assign, bias: 0.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.struct:s .len: 1.. map.attach:s.e .len: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .len: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .len: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .len: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.struct:s .len: 1.. map.attach:s.e .len: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .len: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .len: 8.." 1 "omplower" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-2.c b/gcc/testsuite/c-c++-common/goacc/mdc-2.c new file mode 100644 index 00000000000..ebfb99d4caf --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/mdc-2.c @@ -0,0 +1,62 @@ +/* Test OpenACC's support for manual deep copy, including the attach + and detach clauses. */ + +void +t1 () +{ + struct foo { + int *a, *b, c, d, *e; + } s; + + int *a, *z, scalar, **y; + +#pragma acc enter data copyin(s) detach(z) /* { dg-error ".detach. is not valid for" } */ + { +#pragma acc data copy(s.a[0:10]) copy(z[0:10]) + { + s.e = z; +#pragma acc parallel loop attach(s.e) detach(s.b) /* { dg-error ".detach. is not valid for" } */ + for (int i = 0; i < 10; i++) + s.a[i] = s.e[i]; + + a = s.e; +#pragma acc enter data attach(a) detach(s.c) /* { dg-error ".detach. is not valid for" } */ +#pragma acc exit data detach(a) + } + +#pragma acc enter data attach(z[:5]) /* { dg-error "array section in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(z[:5]) /* { dg-error "array section in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(z[1:]) /* { dg-error "array section in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(z[1:]) /* { dg-error "array section in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(z[:]) /* { dg-error "array section in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(z[:]) /* { dg-error "array section in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(z[3]) /* { dg-error "expected pointer in .attach. clause" } */ +#pragma acc exit data detach(z[3]) /* { dg-error "expected pointer in .detach. clause" } */ + +#pragma acc acc enter data attach(s.e) +#pragma acc exit data detach(s.e) attach(z) /* { dg-error ".attach. is not valid for" } */ + +#pragma acc data attach(s.e) + { + } +#pragma acc exit data delete(a) attach(s.a) /* { dg-error ".attach. is not valid for" } */ + +#pragma acc enter data attach(scalar) /* { dg-error "expected pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(scalar) /* { dg-error "expected pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(s) /* { dg-error "expected pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(s) /* { dg-error "expected pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ + } + +#pragma acc enter data attach(y[10]) +#pragma acc exit data detach(y[10]) +} diff --git a/gcc/testsuite/g++.dg/goacc/mdc.C b/gcc/testsuite/g++.dg/goacc/mdc.C new file mode 100644 index 00000000000..fbd43aaaa60 --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/mdc.C @@ -0,0 +1,68 @@ +/* Test OpenACC's support for manual deep copy, including the attach + and detach clauses. */ + +void +t1 () +{ + struct foo { + int *a, *b, c, d, *e; + } s; + + struct foo& rs = s; + + int *a, *z, scalar, **y; + int* const &ra = a; + int* const &rz = z; + int& rscalar = scalar; + int** const &ry = y; + +#pragma acc enter data copyin(rs) detach(rz) /* { dg-error ".detach. is not valid for" } */ + { +#pragma acc data copy(rs.a[0:10]) copy(rz[0:10]) + { + s.e = z; +#pragma acc parallel loop attach(rs.e) detach(rs.b) /* { dg-error ".detach. is not valid for" } */ + for (int i = 0; i < 10; i++) + s.a[i] = s.e[i]; + + a = s.e; +#pragma acc enter data attach(ra) detach(rs.c) /* { dg-error ".detach. is not valid for" } */ +#pragma acc exit data detach(ra) + } + +#pragma acc enter data attach(rz[:5]) /* { dg-error "array section in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(rz[:5]) /* { dg-error "array section in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(rz[1:]) /* { dg-error "array section in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(rz[1:]) /* { dg-error "array section in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(rz[:]) /* { dg-error "array section in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(rz[:]) /* { dg-error "array section in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(rz[3]) /* { dg-error "expected pointer in .attach. clause" } */ +#pragma acc exit data detach(rz[3]) /* { dg-error "expected pointer in .detach. clause" } */ + +#pragma acc acc enter data attach(rs.e) +#pragma acc exit data detach(rs.e) attach(rz) /* { dg-error ".attach. is not valid for" } */ + +#pragma acc data attach(rs.e) + { + } +#pragma acc exit data delete(ra) attach(rs.a) /* { dg-error ".attach. is not valid for" } */ + +#pragma acc enter data attach(rscalar) /* { dg-error "expected pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(rscalar) /* { dg-error "expected pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(rs) /* { dg-error "expected pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(rs) /* { dg-error "expected pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ + } + +#pragma acc enter data attach(ry[10]) +#pragma acc exit data detach(ry[10]) +} -- 2.17.2