From ee58b02f1160548513a3d00a04acbad93c423b3f Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Fri, 12 May 2017 10:42:31 +0200 Subject: [PATCH] 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. From-SVN: r247953 --- 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(+) create mode 100644 gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c create mode 100644 gcc/testsuite/c-c++-common/goacc/classify-kernels.c create mode 100644 gcc/testsuite/c-c++-common/goacc/classify-parallel.c create mode 100644 gcc/testsuite/c-c++-common/goacc/classify-routine.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 3553c993dde..5ed40a5e2e4 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,14 @@ +2017-05-12 Thomas Schwinge + + * 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 * lib/gcc-dg.exp (schedule-cleanups): Add lang dump capability. diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c new file mode 100644 index 00000000000..a76351c2b8f --- /dev/null +++ b/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 a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c new file mode 100644 index 00000000000..199a73e2a3d --- /dev/null +++ b/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 a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c new file mode 100644 index 00000000000..9d48c1bf9d6 --- /dev/null +++ b/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 a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c new file mode 100644 index 00000000000..72b02c21028 --- /dev/null +++ b/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 a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 new file mode 100644 index 00000000000..fd46d0d506c --- /dev/null +++ b/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 a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 new file mode 100644 index 00000000000..053d27c2a72 --- /dev/null +++ b/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 a/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 new file mode 100644 index 00000000000..087ff48f94f --- /dev/null +++ b/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 a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 new file mode 100644 index 00000000000..319d7678079 --- /dev/null +++ b/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" } } -- 2.30.2