Test cases to check OpenACC offloaded function's attributes and classification

[Thomas Schwinge <thomas at codesourcery dot com>]

Hi!

OK for trunk?

commit 8200af082db5438be18bc60f721fcf21641c0d86
Author: Thomas Schwinge <thomas@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