From 1cf37fc86b0c0f9a36e7150549e00db6f2e6304b Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Fri, 12 May 2017 08:43:45 +0000 Subject: 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: https://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(+) 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.gomp b/gcc/testsuite/ChangeLog.gomp index 952b1010482..b5dd1a45e4f 100644 --- a/gcc/testsuite/ChangeLog.gomp +++ b/gcc/testsuite/ChangeLog.gomp @@ -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-09 Cesar Philippidis * c-c++-common/goacc/update-if_present-1.c: Update test case. 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..d37fb4aea23 --- /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 \\(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 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..dd71a8474d3 --- /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 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" } } -- cgit v1.2.3