diff mbox series

[4/4,OpenACC] Attach / Detach compiler tests

Message ID dfbd80db-af17-9fa7-9a06-d30a92d0bc22@codesourcery.com
State New
Headers show
Series [1/4,OpenACC] Attach / Detach generic infrastructure | expand

Commit Message

Cesar Philippidis Oct. 26, 2018, 4:42 a.m. UTC
This patch introduces a couple of compiler tests for the OpenACC
attach and detach clauses.

Is this OK for trunk after the other patches get approved?

Thanks,
Cesar
diff mbox series

Patch

2018-XX-YY  Cesar Philippidis  <cesar@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.


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.1--