Hi! Ping.
On Wed, 27 Jul 2016 10:59:02 +0200, I wrote:
> 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
signature.asc
Description: PGP signature
