| Message ID | 87zip3jw2x.fsf@hertz.schwinge.homeip.net |
|---|---|
| State | New |
| Headers | show |
Hi! Ping. On Wed, 27 Jul 2016 10:59:02 +0200, I wrote: > 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
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" } }
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(+) Grüße Thomas