Hi!
Ping.
On Thu, 4 Aug 2016 16:06:10 +0200, I wrote:
> Ping.
>
> On Wed, 27 Jul 2016 10:59:02 +0200, I wrote:
> > OK for trunk?
(In the mean time, I also added some more testing.)
commit b7d61270dfc581a6ea130f7a4fa7506a0a5762d8
Author: Thomas Schwinge <tho...@codesourcery.com>
Date: Mon May 8 18:22:50 2017 +0200
Test cases to check OpenACC offloaded function's attributes and
classification
gcc/testsuite/
* c-c++-common/goacc/classify-kernels-unparallelized.c: New file.
* c-c++-common/goacc/classify-kernels.c: Likewise.
* c-c++-common/goacc/classify-parallel.c: Likewise.
* c-c++-common/goacc/classify-routine.c: Likewise.
* gfortran.dg/goacc/classify-kernels-unparallelized.f95: Likewise.
* gfortran.dg/goacc/classify-kernels.f95: Likewise.
* gfortran.dg/goacc/classify-parallel.f95: Likewise.
* gfortran.dg/goacc/classify-routine.f95: Likewise.
---
.../goacc/classify-kernels-unparallelized.c | 39 ++++++++++++++++++++
.../c-c++-common/goacc/classify-kernels.c | 35 ++++++++++++++++++
.../c-c++-common/goacc/classify-parallel.c | 28 +++++++++++++++
.../c-c++-common/goacc/classify-routine.c | 30 ++++++++++++++++
.../goacc/classify-kernels-unparallelized.f95 | 41 ++++++++++++++++++++++
.../gfortran.dg/goacc/classify-kernels.f95 | 37 +++++++++++++++++++
.../gfortran.dg/goacc/classify-parallel.f95 | 30 ++++++++++++++++
.../gfortran.dg/goacc/classify-routine.f95 | 29 +++++++++++++++
8 files changed, 269 insertions(+)
diff --git gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c
gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c
new file mode 100644
index 0000000..a76351c
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c
@@ -0,0 +1,39 @@
+/* Check offloaded function's attributes and classification for unparallelized
+ 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
+
+extern unsigned int *__restrict a;
+extern unsigned int *__restrict b;
+extern unsigned int *__restrict c;
+
+/* An "extern"al mapping of loop iterations/array indices makes the loop
+ unparallelizable. */
+extern unsigned int f (unsigned int);
+
+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[f (i)] + b[f (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 construct is analyzed, and that it
+ can't be parallelized.
+ { dg-final { scan-tree-dump-times "FAILED:" 1 "parloops1" } }
+ { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function
\\(, , \\), omp target entrypoint\\)\\)" 1 "parloops1" } }
+ { dg-final { scan-tree-dump-not "SUCCESS: may be parallelized" "parloops1"
} } */
+
+/* Check the offloaded function's classification and compute dimensions (will
+ always be 1 x 1 x 1 for non-offloading 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" } }
+ { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function
\\(1, 1, 1\\), omp target entrypoint\\)\\)" 1 "oaccdevlow" } } */
diff --git gcc/testsuite/c-c++-common/goacc/classify-kernels.c
gcc/testsuite/c-c++-common/goacc/classify-kernels.c
new file mode 100644
index 0000000..199a73e
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/classify-kernels.c
@@ -0,0 +1,35 @@
+/* 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
+
+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 construct 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)__attribute__\\(\\(oacc function
\\(0, , \\), omp target entrypoint\\)\\)" 1 "parloops1" } }
+ { dg-final { scan-tree-dump-not "FAILED:" "parloops1" } } */
+
+/* Check the offloaded function's classification and compute dimensions (will
+ always be 1 x 1 x 1 for non-offloading 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" } }
+ { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function
\\(1, 1, 1\\), omp target entrypoint\\)\\)" 1 "oaccdevlow" } } */
diff --git gcc/testsuite/c-c++-common/goacc/classify-parallel.c
gcc/testsuite/c-c++-common/goacc/classify-parallel.c
new file mode 100644
index 0000000..9d48c1b
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/classify-parallel.c
@@ -0,0 +1,28 @@
+/* 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
+
+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 x 1 x 1 for non-offloading 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" } }
+ { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function
\\(1, 1, 1\\), omp target entrypoint\\)\\)" 1 "oaccdevlow" } } */
diff --git gcc/testsuite/c-c++-common/goacc/classify-routine.c
gcc/testsuite/c-c++-common/goacc/classify-routine.c
new file mode 100644
index 0000000..72b02c2
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/classify-routine.c
@@ -0,0 +1,30 @@
+/* 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
+
+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 x 1 x 1 for non-offloading 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" } }
+ { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function
\\(0 1, 1 1, 1 1\\), omp declare target, oacc function \\(0 1, 1 0, 1
0\\)\\)\\)" 1 "oaccdevlow" } } */
diff --git gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
new file mode 100644
index 0000000..fd46d0d
--- /dev/null
+++ gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
@@ -0,0 +1,41 @@
+! Check offloaded function's attributes and classification for unparallelized
+! 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
+
+ ! An "external" mapping of loop iterations/array indices makes the loop
+ ! unparallelizable.
+ integer, external :: f
+
+ 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(f (i)) + b(f (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 construct is analyzed, and that it
+! can't be parallelized.
+! { dg-final { scan-tree-dump-times "FAILED:" 1 "parloops1" } }
+! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(,
, \\), omp target entrypoint\\)\\)" 1 "parloops1" } }
+! { dg-final { scan-tree-dump-not "SUCCESS: may be parallelized" "parloops1" }
}
+
+! Check the offloaded function's classification and compute dimensions (will
+! always be 1 x 1 x 1 for non-offloading 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" } }
+! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function
\\(1, 1, 1\\), omp target entrypoint\\)\\)" 1 "oaccdevlow" } }
diff --git gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
new file mode 100644
index 0000000..053d27c
--- /dev/null
+++ gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
@@ -0,0 +1,37 @@
+! 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 construct 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)__attribute__\\(\\(oacc function
\\(0, , \\), omp target entrypoint\\)\\)" 1 "parloops1" } }
+! { dg-final { scan-tree-dump-not "FAILED:" "parloops1" } }
+
+! Check the offloaded function's classification and compute dimensions (will
+! always be 1 x 1 x 1 for non-offloading 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" } }
+! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function
\\(1, 1, 1\\), omp target entrypoint\\)\\)" 1 "oaccdevlow" } }
diff --git gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95
gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95
new file mode 100644
index 0000000..087ff48
--- /dev/null
+++ gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95
@@ -0,0 +1,30 @@
+! 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 x 1 x 1 for non-offloading 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" } }
+! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function
\\(1, 1, 1\\), omp target entrypoint\\)\\)" 1 "oaccdevlow" } }
diff --git gcc/testsuite/gfortran.dg/goacc/classify-routine.f95
gcc/testsuite/gfortran.dg/goacc/classify-routine.f95
new file mode 100644
index 0000000..319d767
--- /dev/null
+++ gcc/testsuite/gfortran.dg/goacc/classify-routine.f95
@@ -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" }
+
+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 x 1 x 1 for non-offloading 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" } }
+! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0
1, 1 1, 1 1\\), omp declare target, oacc function \\(0 0, 1 0, 1 0\\)\\)\\)" 1
"oaccdevlow" } }
Grüße
Thomas
