From: Tom de Vries Date: Wed, 19 Dec 2018 14:20:44 +0000 (+0000) Subject: [nvptx, libgomp] Move rtl-dump test-cases to libgomp X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=49188cd1f2deb943e4047dbffba7a333875d6479;p=gcc.git [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 * 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 --- diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 8f1353465c3..31aff7d80c1 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,9 @@ +2018-12-19 Tom de Vries + + * gcc.dg/goacc/nvptx-merged-loop.c: Move to + libgomp/testsuite/libgomp.oacc-c-c++-common. + * gcc.dg/goacc/nvptx-sese-1.c: Same. + 2018-12-19 Tom de Vries * lib/scanoffloadrtl.exp: New file. diff --git a/gcc/testsuite/gcc.dg/goacc/nvptx-merged-loop.c b/gcc/testsuite/gcc.dg/goacc/nvptx-merged-loop.c deleted file mode 100644 index 3ff537c1d97..00000000000 --- a/gcc/testsuite/gcc.dg/goacc/nvptx-merged-loop.c +++ /dev/null @@ -1,30 +0,0 @@ -/* { dg-do link } */ -/* { dg-require-effective-target offload_nvptx } */ -/* { dg-options "-fopenacc -O2 -foffload=-fdump-rtl-mach\\ -dumpbase\\ nvptx-merged-loop.c\\ -Wa,--no-verify" } */ - -#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-rtl-dump "Merging loop .* into " "mach" } } */ diff --git a/gcc/testsuite/gcc.dg/goacc/nvptx-sese-1.c b/gcc/testsuite/gcc.dg/goacc/nvptx-sese-1.c deleted file mode 100644 index 7e67fe78f06..00000000000 --- a/gcc/testsuite/gcc.dg/goacc/nvptx-sese-1.c +++ /dev/null @@ -1,35 +0,0 @@ -/* { dg-do link } */ -/* { dg-require-effective-target offload_nvptx } */ -/* { dg-options "-fopenacc -O2 -foffload=-fdump-rtl-mach\\ -dumpbase\\ nvptx-sese-1.c\\ -Wa,--no-verify" } */ - -#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-rtl-dump "SESE regions:.* \[0-9\]+{\[0-9\]+->\[0-9\]+(\\.\[0-9\]+)+}" "mach" } } */ diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 084f174513b..ad0abb8a95e 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,10 @@ +2018-12-19 Tom de Vries + + * 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 Chung-Lin Tang diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index c41b3e6dc18..04738a9ce82 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 00000000000..8a2117e1624 --- /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 00000000000..9583265c775 --- /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" } } */