Hi!
On Wed, 10 May 2017 17:49:48 +0200, Jakub Jelinek <ja...@redhat.com> wrote:
> On Mon, May 08, 2017 at 07:02:15PM +0200, Thomas Schwinge wrote:
> > On Thu, 4 Aug 2016 16:06:10 +0200, I wrote:
> > > On Wed, 27 Jul 2016 10:59:02 +0200, I wrote:
> > > > OK for trunk?
> >
> > (In the mean time, I also added some more testing.)
> > Test cases to check OpenACC offloaded function's attributes and
> > classification
> Dunno if it isn't too fragile, but if you are willing to maintain it, ok.
Sure. (Why would I propose it otherwise?)
Committed to trunk in r247953:
commit 692b887e5afc026d8217f0654896f6777edbf7a7
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri May 12 08:42:31 2017 +0000
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.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@247953
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/testsuite/ChangeLog | 11 ++++++
.../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 +++++++++++++++
9 files changed, 280 insertions(+)
diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog
index 3553c99..5ed40a5 100644
--- gcc/testsuite/ChangeLog
+++ gcc/testsuite/ChangeLog
@@ -1,3 +1,14 @@
+2017-05-12 Thomas Schwinge <tho...@codesourcery.com>
+
+ * 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.
+
2017-05-11 Nathan Sidwell <nat...@acm.org>
* lib/gcc-dg.exp (schedule-cleanups): Add lang dump capability.
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" } }
Committed to gomp-4_0-branch in r247954:
commit 5ed11508df8ec3803667636117834b85005f5990
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri May 12 08:43:45 2017 +0000
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.
trunk r247953
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@247954
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/testsuite/ChangeLog.gomp | 11 ++++++
.../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 +++++++++++++++
9 files changed, 280 insertions(+)
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 952b101..b5dd1a4 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,14 @@
+2017-05-12 Thomas Schwinge <tho...@codesourcery.com>
+
+ * 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.
+
2017-05-09 Cesar Philippidis <ce...@codesourcery.com>
* c-c++-common/goacc/update-if_present-1.c: Update test case.
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..d37fb4a
--- /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 \\(worker\\), 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 \\(worker\\), 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..dd71a84
--- /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 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" } }
Grüße
Thomas
