Hi!
OK for trunk?
commit 8200af082db5438be18bc60f721fcf21641c0d86
Author: Thomas Schwinge <tho...@codesourcery.com>
Date: Tue Jul 26 17:18:21 2016 +0200
Test cases to check OpenACC offloaded function's attributes and
classification
gcc/testsuite/
* c-c++-common/goacc/oaccdevlow-kernels.c: New file.
* c-c++-common/goacc/oaccdevlow-parallel.c: Likewise.
* c-c++-common/goacc/oaccdevlow-routine.c: Likewise.
* gfortran.dg/goacc/oaccdevlow-kernels.f95: Likewise.
* gfortran.dg/goacc/oaccdevlow-parallel.f95: Likewise.
* gfortran.dg/goacc/oaccdevlow-routine.f95: Likewise.
---
.../c-c++-common/goacc/oaccdevlow-kernels.c | 34 ++++++++++++++++++++
.../c-c++-common/goacc/oaccdevlow-parallel.c | 27 ++++++++++++++++
.../c-c++-common/goacc/oaccdevlow-routine.c | 29 +++++++++++++++++
.../gfortran.dg/goacc/oaccdevlow-kernels.f95 | 36 ++++++++++++++++++++++
.../gfortran.dg/goacc/oaccdevlow-parallel.f95 | 29 +++++++++++++++++
.../gfortran.dg/goacc/oaccdevlow-routine.f95 | 28 +++++++++++++++++
6 files changed, 183 insertions(+)
diff --git gcc/testsuite/c-c++-common/goacc/oaccdevlow-kernels.c
gcc/testsuite/c-c++-common/goacc/oaccdevlow-kernels.c
new file mode 100644
index 0000000..14d650a
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/oaccdevlow-kernels.c
@@ -0,0 +1,34 @@
+/* Check offloaded function's attributes and classification for OpenACC
+ kernels. */
+
+/* { dg-additional-options "-O2" }
+ { dg-additional-options "-fdump-tree-ompexp" }
+ { dg-additional-options "-fdump-tree-parloops1-all" }
+ { dg-additional-options "-fdump-tree-oaccdevlow" } */
+
+#define N (1024 * 512)
+
+extern unsigned int *__restrict a;
+extern unsigned int *__restrict b;
+extern unsigned int *__restrict c;
+
+void KERNELS ()
+{
+#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
+ for (unsigned int i = 0; i < N; i++)
+ c[i] = a[i] + b[i];
+}
+
+/* Check the offloaded function's attributes.
+ { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp target
entrypoint\\)\\)" 1 "ompexp" } } */
+
+/* Check that exactly one OpenACC kernels loop is analyzed, and that it can be
+ parallelized.
+ { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1
"parloops1" } }
+ { dg-final { scan-tree-dump-times "(?n)oacc function \\(0," 1 "parloops1" }
}
+ { dg-final { scan-tree-dump-not "FAILED:" "parloops1" } } */
+
+/* Check the offloaded function's classification and compute dimensions (will
+ always be [1, 1, 1] for target compilation).
+ { dg-final { scan-tree-dump-times "(?n)Function is kernels offload" 1
"oaccdevlow" } }
+ { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]"
1 "oaccdevlow" } } */
diff --git gcc/testsuite/c-c++-common/goacc/oaccdevlow-parallel.c
gcc/testsuite/c-c++-common/goacc/oaccdevlow-parallel.c
new file mode 100644
index 0000000..63c372a
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/oaccdevlow-parallel.c
@@ -0,0 +1,27 @@
+/* Check offloaded function's attributes and classification for OpenACC
+ parallel. */
+
+/* { dg-additional-options "-O2" }
+ { dg-additional-options "-fdump-tree-ompexp" }
+ { dg-additional-options "-fdump-tree-oaccdevlow" } */
+
+#define N (1024 * 512)
+
+extern unsigned int *__restrict a;
+extern unsigned int *__restrict b;
+extern unsigned int *__restrict c;
+
+void PARALLEL ()
+{
+#pragma acc parallel loop copyin (a[0:N], b[0:N]) copyout (c[0:N])
+ for (unsigned int i = 0; i < N; i++)
+ c[i] = a[i] + b[i];
+}
+
+/* Check the offloaded function's attributes.
+ { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp target
entrypoint\\)\\)" 1 "ompexp" } } */
+
+/* Check the offloaded function's classification and compute dimensions (will
+ always be [1, 1, 1] for target compilation).
+ { dg-final { scan-tree-dump-times "(?n)Function is parallel offload" 1
"oaccdevlow" } }
+ { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]"
1 "oaccdevlow" } } */
diff --git gcc/testsuite/c-c++-common/goacc/oaccdevlow-routine.c
gcc/testsuite/c-c++-common/goacc/oaccdevlow-routine.c
new file mode 100644
index 0000000..fa2eae7
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/oaccdevlow-routine.c
@@ -0,0 +1,29 @@
+/* Check offloaded function's attributes and classification for OpenACC
+ routine. */
+
+/* { dg-additional-options "-O2" }
+ { dg-additional-options "-fdump-tree-ompexp" }
+ { dg-additional-options "-fdump-tree-oaccdevlow" } */
+
+#define N (1024 * 512)
+
+extern unsigned int *__restrict a;
+extern unsigned int *__restrict b;
+extern unsigned int *__restrict c;
+#pragma acc declare copyin (a, b) create (c)
+
+#pragma acc routine worker
+void ROUTINE ()
+{
+#pragma acc loop
+ for (unsigned int i = 0; i < N; i++)
+ c[i] = a[i] + b[i];
+}
+
+/* Check the offloaded function's attributes.
+ { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare
target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */
+
+/* Check the offloaded function's classification and compute dimensions (will
+ always be [1, 1, 1] for target compilation).
+ { dg-final { scan-tree-dump-times "(?n)Function is routine level 1" 1
"oaccdevlow" } }
+ { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]"
1 "oaccdevlow" } } */
diff --git gcc/testsuite/gfortran.dg/goacc/oaccdevlow-kernels.f95
gcc/testsuite/gfortran.dg/goacc/oaccdevlow-kernels.f95
new file mode 100644
index 0000000..8ee641e
--- /dev/null
+++ gcc/testsuite/gfortran.dg/goacc/oaccdevlow-kernels.f95
@@ -0,0 +1,36 @@
+! Check offloaded function's attributes and classification for OpenACC
+! kernels.
+
+! { dg-additional-options "-O2" }
+! { dg-additional-options "-fdump-tree-ompexp" }
+! { dg-additional-options "-fdump-tree-parloops1-all" }
+! { dg-additional-options "-fdump-tree-oaccdevlow" }
+
+program main
+ implicit none
+ integer, parameter :: n = 1024
+ integer, dimension (0:n-1) :: a, b, c
+ integer :: i
+
+ call setup(a, b)
+
+ !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1))
+ do i = 0, n - 1
+ c(i) = a(i) + b(i)
+ end do
+ !$acc end kernels
+end program main
+
+! Check the offloaded function's attributes.
+! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp target
entrypoint\\)\\)" 1 "ompexp" } }
+
+! Check that exactly one OpenACC kernels loop is analyzed, and that it can be
+! parallelized.
+! { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1
"parloops1" } }
+! { dg-final { scan-tree-dump-times "(?n)oacc function \\(0," 1 "parloops1" } }
+! { dg-final { scan-tree-dump-not "FAILED:" "parloops1" } }
+
+! Check the offloaded function's classification and compute dimensions (will
+! always be [1, 1, 1] for target compilation).
+! { dg-final { scan-tree-dump-times "(?n)Function is kernels offload" 1
"oaccdevlow" } }
+! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1
"oaccdevlow" } }
diff --git gcc/testsuite/gfortran.dg/goacc/oaccdevlow-parallel.f95
gcc/testsuite/gfortran.dg/goacc/oaccdevlow-parallel.f95
new file mode 100644
index 0000000..0975eb8
--- /dev/null
+++ gcc/testsuite/gfortran.dg/goacc/oaccdevlow-parallel.f95
@@ -0,0 +1,29 @@
+! Check offloaded function's attributes and classification for OpenACC
+! parallel.
+
+! { dg-additional-options "-O2" }
+! { dg-additional-options "-fdump-tree-ompexp" }
+! { dg-additional-options "-fdump-tree-oaccdevlow" }
+
+program main
+ implicit none
+ integer, parameter :: n = 1024
+ integer, dimension (0:n-1) :: a, b, c
+ integer :: i
+
+ call setup(a, b)
+
+ !$acc parallel loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1))
+ do i = 0, n - 1
+ c(i) = a(i) + b(i)
+ end do
+ !$acc end parallel loop
+end program main
+
+! Check the offloaded function's attributes.
+! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp target
entrypoint\\)\\)" 1 "ompexp" } }
+
+! Check the offloaded function's classification and compute dimensions (will
+! always be [1, 1, 1] for target compilation).
+! { dg-final { scan-tree-dump-times "(?n)Function is parallel offload" 1
"oaccdevlow" } }
+! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1
"oaccdevlow" } }
diff --git gcc/testsuite/gfortran.dg/goacc/oaccdevlow-routine.f95
gcc/testsuite/gfortran.dg/goacc/oaccdevlow-routine.f95
new file mode 100644
index 0000000..a68b5eb
--- /dev/null
+++ gcc/testsuite/gfortran.dg/goacc/oaccdevlow-routine.f95
@@ -0,0 +1,28 @@
+! Check offloaded function's attributes and classification for OpenACC
+! routine.
+
+! { dg-additional-options "-O2" }
+! { dg-additional-options "-fdump-tree-ompexp" }
+! { dg-additional-options "-fdump-tree-oaccdevlow" }
+
+subroutine ROUTINE
+ !$acc routine worker
+ integer, parameter :: n = 1024
+ integer, dimension (0:n-1) :: a, b, c
+ integer :: i
+
+ call setup(a, b)
+
+ !$acc loop
+ do i = 0, n - 1
+ c(i) = a(i) + b(i)
+ end do
+end subroutine ROUTINE
+
+! Check the offloaded function's attributes.
+! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare
target, oacc function \\(0 0, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } }
+
+! Check the offloaded function's classification and compute dimensions (will
+! always be [1, 1, 1] for target compilation).
+! { dg-final { scan-tree-dump-times "(?n)Function is routine level 1" 1
"oaccdevlow" } }
+! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1
"oaccdevlow" } }
Grüße
Thomas
