diff options
author | Tom de Vries <tdevries@suse.de> | 2018-12-19 14:20:44 +0000 |
---|---|---|
committer | Tom de Vries <vries@gcc.gnu.org> | 2018-12-19 14:20:44 +0000 |
commit | 49188cd1f2deb943e4047dbffba7a333875d6479 (patch) | |
tree | e2a4813813fd016575e3cbc5e7f561a7c933b117 /libgomp | |
parent | 9b09e453c6884f2ea391dc70c188c159f0883150 (diff) | |
download | gcc-49188cd1f2deb943e4047dbffba7a333875d6479.zip gcc-49188cd1f2deb943e4047dbffba7a333875d6479.tar.gz gcc-49188cd1f2deb943e4047dbffba7a333875d6479.tar.bz2 |
[nvptx, libgomp] Move rtl-dump test-cases to libgomp
The goacc.exp test-cases nvptx-merged-loop.c and nvptx-sese-1.c are failing
during linking due to missing libgomp.spec.
Move them to the libgomp testsuite.
Build and reg-tested on x86_64 with nvptx accelerator.
2018-12-19 Tom de Vries <tdevries@suse.de>
* gcc.dg/goacc/nvptx-merged-loop.c: Move to
libgomp/testsuite/libgomp.oacc-c-c++-common.
* gcc.dg/goacc/nvptx-sese-1.c: Same.
* testsuite/lib/libgomp.exp: Add load_lib of scanoffloadrtl.exp.
* testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c: Move from
gcc/testsuite/gcc.dg/goacc.
* testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c: Same.
From-SVN: r267267
Diffstat (limited to 'libgomp')
-rw-r--r-- | libgomp/ChangeLog | 7 | ||||
-rw-r--r-- | libgomp/testsuite/lib/libgomp.exp | 1 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c | 30 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c | 35 |
4 files changed, 73 insertions, 0 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 084f174..ad0abb8 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,10 @@ +2018-12-19 Tom de Vries <tdevries@suse.de> + + * testsuite/lib/libgomp.exp: Add load_lib of scanoffloadrtl.exp. + * testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c: Move from + gcc/testsuite/gcc.dg/goacc. + * testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c: Same. + 2018-12-14 Thomas Schwinge <thomas@codesourcery.com> Chung-Lin Tang <cltang@codesourcery.com> diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index c41b3e6..04738a9 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -31,6 +31,7 @@ load_gcc_lib scanrtl.exp load_gcc_lib scantree.exp load_gcc_lib scanltranstree.exp load_gcc_lib scanoffloadtree.exp +load_gcc_lib scanoffloadrtl.exp load_gcc_lib scanipa.exp load_gcc_lib scanwpaipa.exp load_gcc_lib timeout-dg.exp diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c new file mode 100644 index 0000000..8a2117e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c @@ -0,0 +1,30 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ +/* { dg-options "-foffload=-fdump-rtl-mach" } */ +/* { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ + +#define N (32*32*32+17) +void __attribute__ ((noinline)) Foo (int *ary) +{ + int ix; + +#pragma acc parallel num_workers(32) vector_length(32) copyout(ary[0:N]) + { + /* Loop partitioning should be merged. */ +#pragma acc loop worker vector + for (unsigned ix = 0; ix < N; ix++) + { + ary[ix] = ix; + } + } +} + +int main () +{ + int ary[N]; + + Foo (ary); + + return 0; +} + +/* { dg-final { scan-offload-rtl-dump "Merging loop .* into " "mach" } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c new file mode 100644 index 0000000..9583265 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c @@ -0,0 +1,35 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ +/* { dg-options "-foffload=-fdump-rtl-mach" } */ +/* { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ + +#pragma acc routine seq +int __attribute__((noinline)) foo (int x) +{ + return x & 2; +} + +int main () +{ + int r = 0; + +#pragma acc parallel copy(r) vector_length(32) + { +#pragma acc loop vector reduction (+:r) + for (int i = 00; i < 40; i++) + r += i; + + /* This piece is a multi-block SESE region */ + if (foo (r)) + r *= 2; + + if (r & 1) /* to here. */ +#pragma acc loop vector reduction (+:r) + for (int i = 00; i < 40; i++) + r += i; + } + + return 0; +} + +/* Match {N->N(.N)+} */ +/* { dg-final { scan-offload-rtl-dump "SESE regions:.* \[0-9\]+{\[0-9\]+->\[0-9\]+(\\.\[0-9\]+)+}" "mach" } } */ |